从给定的二维偏移中知道CUDA中的块ID

Question

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

例如，如果我有一个由64个元素组成的1D数组，并且我用16x1个线程（总共4个块）配置内核，则每个线程都可以使用访问一个位置

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

并且我可以很容易地从idx中获得给定索引值的blockIdx.x作为

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

如何知道给定idx的块ID.x和ID.y？即索引34和35在块（0,1）中，索引36在块（1,1）中。因此，如果块（0，1）中的线程读取索引35的值，则该线程将知道该值在其块内，并将从共享内存中读取该值。索引35的值将被存储在块（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和block Idx.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）。