从给定的 2D 偏移量了解 CUDA 中的块 ID

Question

我试图根据 CUDA 中给定的偏移量计算 blockIdx.x 和 blockIdx.y，但我完全无法思考。这个想法是在可能的情况下从共享内存中读取数据，在其他情况下从全局内存中读取数据。

例如，如果我有一个包含 64 个元素的一维数组，并且我配置了一个具有 16x1 线程（总共 4 个块）的内核，每个线程可以使用以下命令访问一个位置：

int idx = blockDim.x*blockIdx.x + threadIdx.x

并且我可以轻松获取给定索引值的 blockIdx.x从 idx 来看，

int blockNumber = idx / blockDim.x; 

但在具有 8x8 元素和 4x4 线程内核配置（总共 2x2 块）的 2D 场景中，每个线程使用以下方式访问一个位置：

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int pitch = blockDim.x * gridDim.x;
int idx = x + y * pitch;

int sharedMemIndex = threadIdx.x+threadIdx.y+BLOCK_DIM_X;
__shared_block[sharedMemIndex] = fromGlobalMemory[idx];
__syncthreads();

// ... some operations

int unknow_index = __shared_block[sharedMemIndex];

if ( unknow_index within this block? )
    // ... read from shared memory
else
    // ... read from global memory

我如何知道块 ID.x 和ID.y 在给定的 idx 处？即索引34和35在块(0, 1)中，索引36在块(1, 1)中。因此，如果块 (0, 1) 中的线程读取索引 35 的值，则该线程将知道该值在其块内，并将从共享内存中读取它。索引 35 的值将存储在该块的共享内存的位置 11 (0. 1) 中。

提前致谢！

Answer 1

在实践中，我真的想不出为什么需要这样做的充分理由，但是您可以这样计算结果，对于任意索引值 idx（假设列有序索引）：

int pitch = blockDim.x * gridDim.x;
int tidy = idx / pitch; // div(idx,pitch)
int tidx = idx - (pitch * tidy); // mod(idx,pitch)
int bidx = idx / blockDim.x;
int bidy = idy / blockDim.y;

应该给出 bidx 和 bidy 中索引的块坐标。

Answer 2

无需对 Idx 应用数学来找出 X 和 Y 块，也无需从 Idx 向后查找块索引。对于每个线程 (Idx)，您只需调用 blockIdx.x 和 blockIdx.y 即可找到 Y 和 X 块。

在内核中的任何点：

int x = blockIdx.x // will give you X block Index at that particular thread
int y = blockIdx.y // will give you Y block Index at that particular thread. 

更新：
如果您执着于反向操作，则需要知道节距和块尺寸的值

   int currentRow = idx/pitch;
   int currentCol = idx%pitch;

   int block_idx_x = currentCol/blockDim.x;
   int block_idx_y = currentRow/blockDim.y;

Answer 3

您正在执行不必要的计算。

idx / blockDim.x
-->(blockDim.x * blockIdx.x + threadIdx.x)/blockDim.x
-->(blockIdx.x  + threadIdx.x/blockDim.x)
--> blockIdx.x + 0 (threadIdx.x always less than blockDim.x)

您可以只使用 blockIdx.x 而不是复杂的计算。对于 2D 网格（blockIdx.x 和 blockIdx.y）也是如此。