CUDA：在使用 cudaMallocPitch 分配的二维数组中查找数组索引

Question

我一直在用 cuda 开发一款生活游戏。我想找到每个元素的数组索引，以便我可以计算该元素的邻居，以便我可以将新值写入该数组。我发现的所有关于此的事情都在使用指向行的指针，我只是无法弄清楚这究竟是如何转换为索引的。为了更好地理解我的意思，我必须遵循代码（一些 sn-ps）：

#define THREADSPERBLOCK 1024
lifeState *d_gameStateInitial;
size_t d_gameStateInitial_pitch;
int sizeX = 100;
int sizeY = 100;

int numBlocks = ((sizeX * sizeY) % THREADSPERBLOCK) + 1;
int numThreadsPerBlock;

if(numBlocks == 1)
{
    numThreadsPerBlock = sizeX * sizeY;
}
else
{
    numThreadsPerBlock = THREADSPERBLOCK;
}

cudaMallocPitch((void **)&d_gameStateInitial, &d_gameStateInitial_pitch, sizeX * sizeof(lifeState), sizeY);

doTheGame<<<numBlocks, numThreadsPerBlock>>>(d_gameStateInitial, d_gameStateInitial_pitch, d_gameStateNew, d_gameStateNew_pitch, sizeX, sizeY);

“lifestate *”只是一个包含死/活枚举的结构。两个数组，初始数组和新数组都以完全相同的方式进行 malloc。在 doTheGame 内核中，我现在想知道如何计算索引，我正在考虑这样的事情，但我认为这是错误的：

__global__ void doTheGame(lifeState *initialArray, size_t initialArrayPitch,
                      lifeState *newArray, size_t newArrayPitch,
                      int sizeX, int sizeY)
{
int initialArrayThreadIndex = (blockIdx.x * initialArrayPitch) + threadIdx.x;
int newArrayThreadIndex = (blockIdx.x * initialArrayPitch) + threadIdx.x;
}

到目前为止，我发现的所有内容都与 cudaMallocPitch 示例基本相同：

  T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;

但我只是看不出它如何准确地转换为块、线程以及 x 和 y。

提前致谢。

Answer 1

假设我有一个像这样的double data[] 数组：

A  B  C  D  x  x  x  x
E  F  G  H  x  x  x  x

然后是data[0] = A、data[1] = B 等

这可能是分配一个间距为 64 字节的 2x4 数组的结果（因此每行中有 4 个额外的 x 条目）。假设上述数组中的每个元素都是一个double 数量。

现在假设我有一个内核，并且我正在启动一个 2x4 线程数组，每个有效元素一个（x 元素无效 - 它们是满足音高要求的额外分配，这是任意选择的这个例子是 64 字节）。

在这个内核中，我可以创建一个 x 和 y 索引，如下所示：

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

然后每个线程可以访问其各自的元素，如下所示。

由于cudaMallocPitch返回的pitch数量是以字节为单位的，我们需要先计算一个行偏移量，使用char指针算法：

double *rowptr = (double *)((char *)data + idy*pitch);

然后我们可以像这样访问该行上的元素：

rowptr[idx] = 0.0;

如果我希望线程访问与其线程索引对应的元素以外的元素，则计算类似。

例如在上面的data数组中将元素G（即元素（1,2））设置为零，我可以这样做：

double *rowptr = (double *)((char *)data + 1*pitch);
rowptr[2] = 0.0;