CUDA. Использование Warp Shuffle Functions в случае если blockDim > warpSize

Разбираюсь с cuda и не могу понять одну вещь. Как мне синхронизировать результаты warp shuffle функции, если размер блока больше размера варпа?

Предположим я хочу просуммировать числа с помощью __shfl_down_sync и размер блока (8,8). Вот функция, которая работает для случаев, когда размер блока не превышает 32.

cudaError_t cudaTestWarpShuffle8x8BlockDim(cudaStream_t stream)
{
    const dim3 blockDim(8, 8);
    const dim3 gridDim(1, 1);
    gpuTestWarpShuffle8x8BlockDim <<<gridDim, blockDim, 0, stream>>>();
    return cudaSuccess;
}

#define FULL_MASK_ 0x0ffffffff
__global__ void gpuTestWarpShuffle8x8BlockDim()
{
    int threadNum = threadIdx.x + threadIdx.y * blockDim.x;
    int used = threadNum < 49 ? 1 : 0;

    for (int offset = blockDim.x * blockDim.y / 2; offset > 0; offset /= 2)
    {
        used += __shfl_down_sync(FULL_MASK_, used, offset);
    }
    if (threadIdx.x == 0 && threadIdx.y == 0)
    {
         printf("CUDA DEBUG [threadNum : %d, sum: %d]\n", threadNum, used);
    }
}

Не очень понимаю, как ее модифицировать чтобы синхронизировать варпы внутри блока и получить нужную сумму? (в данном случае 49)