CUDA 3D 到线性索引映射（Pitch）答案

【问题标题】：CUDA 3D to linear index mapping (Pitch)CUDA 3D 到线性索引映射（Pitch）
【发布时间】：2015-02-27 14:07:35
【问题描述】：

我有 3 维数据，我正在使用 CUDA 处理它。

我正在使用 cudaMallocPitch() 分配内存。

cudaMallocPitch((void **)&test_data, &pitch, sizeof(float)*N*N, N);

在二维版本的代码中，维度为 N*N，我通过以下方式访问特定元素：

i = blockIdx.x*BLOCK_X + threadIdx.x;
j = blockIdx.y*BLOCK_Y + threadIdx.y;
linearIdx = i + j*pitch/sizeof(float);

现在我希望将代码扩展到 3-D。我可以类似地得到 z 索引

k = blockIdx.z*BLOCK_Z + threadIdx.z;

但是我现在如何将这三者结合起来得到线性索引呢？ pitch 到底是什么，我现在如何访问该元素？请评论我为 3-D 数据分配内存的方法是否正确。

谢谢！

【问题讨论】：

我不知道这个功能 (cudaMallocPitch)，但它似乎有一个 3D 版本，我通过以下 cudaMalloPitch cudaMallocPitch 上的第一个谷歌链接找到它：developer.download.nvidia.com/compute/cuda/4_1/rel/toolkit/docs/… 所以我猜你（谁知道）也可以找到它，公式可能是 linearIdx = i + pitch.xsize * ( j + pitch.ysize * k)
亲爱的@user2774555，如果您重新审核答案并接受其中一个答案，或者甚至使用新信息编辑您的问题，如果没有帮助您解决问题，那就太好了。

标签： 3d cuda pitch

【解决方案1】：

在 2D 和 3D 数组的编程指南中有几个很好的示例。它们如下：

二维：

// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,
                width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);

// Device code
__global__ void MyKernel(float* devPtr,
                         size_t pitch, int width, int height)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
            float element = row[c];
        }
    }
}

3D：

// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),
                                    height, depth);
cudaPitchedPtr devPitchedPtr;
cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);

// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,
                         int width, int height, int depth)
{
    char* devPtr = devPitchedPtr.ptr;
    size_t pitch = devPitchedPtr.pitch;
    size_t slicePitch = pitch * height;
    for (int z = 0; z < depth; ++z) {
        char* slice = devPtr + z * slicePitch;
        for (int y = 0; y < height; ++y) {
            float* row = (float*)(slice + y * pitch);
            for (int x = 0; x < width; ++x) {
                float element = row[x];
            }
        }
    }
}

使用 2D 数组很容易看到音高的使用情况。他们将数组指针转换为char* 的原因是，pitch 返回一个字节大小，而不是元素数量（pitch 可能不是元素大小的倍数）。

对于 3D 数组，这只是使用每个 2D 数组的高度进行扩展。这类似于将 3D 结构展开为许多 2D 切片。

【讨论】：

你这里没有使用blockIdx和threadIdx。谢谢你推荐cudaMalloc3D()函数，但是如何合并线程索引和块索引？
根据 3D 结构的大小，有几种不同的方法。您可以为数组的每个 2D 切片分配一个块，然后使用 2D 块结构循环遍历元素块。您可以为每个 2D 切片使用一个块，然后使用一个 1D 块循环遍历每个 2D 切片的行。您是否对某个特定的映射感兴趣？

【解决方案2】：

内存是一维连续的字节空间。 1D、2D 和 3D 访问模式取决于您如何解释数据，以及您如何通过 1D、2D 和 3D 线程块访问它们。

cudaMallocPitch 在设备上分配至少宽度（以字节为单位）* 高度字节的线性内存。函数可以填充分配以确保任何给定行中的相应指针在地址逐行更新时将继续满足合并的对齐要求。间距由 cudaMallocPitch() 在 *pitch 中返回 是分配的宽度（以字节为单位）。 - 内存管理 [CUDA 运行时 API]。

在M (rows) x N (cols) x K (slices)的数据情况下，一个切片的每个像素将在

i = blockIdx.x*BLOCK_X + threadIdx.x;
j = blockIdx.y*BLOCK_Y + threadIdx.y;

linearIdx = i + j*pitch/sizeof(float);

下一个像素切片是M x N 位置分隔。因此，要访问您的数据，您必须正确跳转到下一个切片。也就是说，

i = blockIdx.x*BLOCK_X + threadIdx.x;
j = blockIdx.y*BLOCK_Y + threadIdx.y;
// index for the slice
k = blockIdx.z*BLOCK_Z + threadIdx.z;

// for the sake of simplicity 
int next_row_pitched = pitch/sizeof(float);

linearIdx = i + j*next_row_pitched + k*next_row_pitched*N;

要获得“第三”维度，您必须“跳跃”一个切片的所有像素，即M x N 位置。随着每一行的倾斜，您必须将 M 更改为 cudaMallockPitch 返回的倾斜值。

如果您启动 3D 线程块，上述索引是有效的。您还可以启动 2D 线程块，然后迭代切片数。

CUDA C 编程指南的第 3.2.2 章设备内存有一个代码示例，它分配一个宽×高×深的浮点值 3D 数组，并展示了如何在设备代码中循环遍历数组元素。

【讨论】：