管理 2D CUDA 阵列答案

【问题标题】：Managing a 2D CUDA Array管理 2D CUDA 阵列
【发布时间】：2013-09-02 05:03:06
【问题描述】：

我正在尝试将二维数组传递给内核，以便每个线程都可以访问 index = threadIdx.x + (blockIdx.x * blockDim.x) 但我无法弄清楚如何执行此操作如何将数据复制回来。

size_t pitch;
cudaMallocPitch(&d_array, &pitch, block_size * sizeof(int), num_blocks);
cudaMemset2D(d_array, pitch, 0, block_size * sizeof(int), num_blocks * sizeof(int));
kernel<<<grid_size, block_size>>>(d_array, pitch);
cudaMemcpy2D(h_array, pitch, d_array, pitch, block_size, num_blocks, cudaMemcpyDeviceToHost);
for (num_blocks)
  for(block_size)
    h_array[block][thread] should be 1

__global__ void kernel(int *array, int pitch) {
  int *row = (int*)((char*)array + blockIdx.x * pitch);
  row[threadIdx.x] = 1;
  return;
}

我做错了什么，在这里？

【问题讨论】：

为什么要将数组转换为 (char*) ？这将导致错误的指针算法
这两个问题是这样描述的：stackoverflow.com/questions/1047369/…stackoverflow.com/questions/5029920/…
@LarryPel：不，不会。间距以字节为单位，需要一个指向字节大小类型的指针才能正确执行指针运算。
请提供一个完整的、可编译的程序。不要让我们玩什么是“block_size”，什么是“grid_size”等20个问题。

标签： c++ c arrays cuda 2d

【解决方案1】：

您的 cudaMemset2D 正在访问您之前使用 cudaMallocPitch 分配的更大内存空间此外，您的 cudaMemcpy2D 正在复制该内存的一小部分。

您应该按以下方式使用该功能：

cudaMallocPitch(&d_array, &pitch, block_size * sizeof(int), num_blocks);
cudaMemset2D(d_array, pitch, 0, block_size * sizeof(int), num_blocks) // * sizeof(int)); <- This size is bigger than the previously declared
kernel<<<grid_size, block_size>>>(d_array, pitch);
cudaMemcpy2D(h_array, pitch, d_array, pitch, block_size * sizeof(int) /* you forgot this here */, num_blocks, cudaMemcpyDeviceToHost);

【讨论】：

另外，如果你做了proper cuda error checking 你的cudaMemset2D（至少）会抛出一个错误。
另外，大概你的h_array 没有投球。因此，您在 cudaMemcpy2D 中为其传递的 pitch 参数不正确。你可能应该通过block_size*sizeof(int)（或类似的东西）来获得h_array的音高。

【解决方案2】：

这是一个通过基本测试的完整代码，修复了@hidrargyro 提到的错误：

$ cat t236.cu
#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


__global__ void kernel(int *array, int pitch) {
  int *row = (int*)((char*)array + blockIdx.x * pitch);
  row[threadIdx.x] = 1;
  return;
}

int main(){

int *d_array, *h_array;
int block_size = 256;
int num_blocks = 256;
int grid_size = num_blocks;
h_array=(int *)malloc(block_size*num_blocks*sizeof(int));
if (h_array==0) {printf("malloc fail\n"); return 1;}
cudaMalloc((void **)&d_array, block_size*num_blocks*sizeof(int));
cudaCheckErrors("cudaMalloc fail");

size_t pitch;
cudaMallocPitch(&d_array, &pitch, block_size * sizeof(int), num_blocks);
cudaCheckErrors("cudaMallocPitch fail");
cudaMemset2D(d_array, pitch, 0, block_size * sizeof(int), num_blocks);
cudaCheckErrors("cudaMemset2D fail");
kernel<<<grid_size, block_size>>>(d_array, pitch);
cudaDeviceSynchronize();
cudaCheckErrors("kernel fail");

cudaMemcpy2D(h_array, block_size*sizeof(int), d_array, pitch, block_size*sizeof(int), num_blocks, cudaMemcpyDeviceToHost);
cudaCheckErrors("cudaMemcpy 2D fail");
for (int i = 0; i<num_blocks; i++)
  for(int j = 0; j<block_size; j++)
    if (h_array[i*block_size+j] != 1) {printf("mismatch at i=%d, j=%d, should be 1, was %d\n", i,j,h_array[i*block_size+j]); return 1;}
printf("success\n");
return 0;
}

$ nvcc -arch=sm_20 -o t236 t236.cu
$ ./t236
success
$

如果您打算接受答案，请接受@hidrargyro 给出的答案

【讨论】：

这就像 CUDA_SAFE_CALL()，我总是想知道 while(0) 的奇怪循环