访问用 cudaMallocPitch 分配的内存答案

【问题标题】：accessing memory allocated with cudaMallocPitch访问用 cudaMallocPitch 分配的内存
【发布时间】：2014-02-16 02:23:40
【问题描述】：

我正在尝试在 cuda 中创建一个二维数组并对其进行初始化，但失败得很惨。这是我的初始化内核：

__global__ void initMap(float* map, size_t pitch, int w, int h, int numX, int numY){
int idx=blockIdx.x*blockDim.x+threadIdx.x;
int idy=blockIdx.y*blockDim.y;
int i;
int j;
for (i=0; i<numX; i++){
    for (j=0; j<numY; j++){
        int idMC=idx+i;
        int idMR=threadIdx.y+j;
        if(idMC<w && idy+idMR<h){
             float* row=(float*)(map+idy+idMR*pitch);
             row[idMC]=0.5;
        }
    }
}

__syncthreads();
}

下面是我如何分配数组并在 main 中调用初始化内核：

int width=map_size;
int height=map_size;
float* map;
size_t pitch;
checkCudaErrors(cudaMallocPitch(&map, &pitch, width*sizeof(float), height));
int numT=32;
int numBX=(int)ceil((float)width/numT);
int numBY=(int)ceil((float)height/numT);
dim3 numBlocks(numBX, numBY);
dim3 numThr(numT, numT);
initMap <<<numBlocks, numThr>>> (map, pitch/sizeof(float), width, height, 1, 1);
cudaError_t err=cudaGetLastError();
if (err != cudaSuccess) 
    printf("Error: %s\n", cudaGetErrorString(err));
checkCudaErrors(cudaDeviceSynchronize());

基本上，我要做的是将 2D 数组划分为 32x32 块，并将每个块分配给一个块，然后用 0.5 填充它。但是，当我尝试将阵列写入磁盘或从另一个内核访问其元素时，我可以看到那里有很多 QNAN，所以我想我并没有填满整个阵列，而是遗漏了一些地方。

【问题讨论】：

没关系，我只是在初始化时忘记在 idy+idMR 周围加上括号。
这仍然不正确。虽然它可能会起作用，但它只是偶然发生的。不能保证音高是分配字大小的整数倍。 cudaMallocPitch 的文档包括内核代码中音高的正确使用。
它是转换为 (char*) 的那个，对吧？我没有使用它，因为我不太清楚那里正在做什么。你介意给我解释一下吗？谢谢！

标签： cuda gpu gpgpu

【解决方案1】：

这样的事情应该可以工作：

__global__ void initMap(float* map, size_t pitch, int w, int h, int numX, int numY){
  int col = threadIdx.x + (blockDim.x * blockIdx.x);
  int row = threadIdx.y + (blockDim.y * blockIdx.y);

  if ((row < h) && (col < w)){
    float *myrow = (float *)(((char *)map)+(row*pitch));
    myrow[col] = 0.5f;
}

这假定pitch 参数将作为字节的数量传递，这是cudaMallocPitch 设置它的方式。所以我们要做的第一件事就是将map 转换为char 指针。然后我们根据选定的行为其添加适当的倾斜偏移，并将结果指针转换为float 指针。然后我们使用float 指针索引到选定的行。

因此你的内核调用应该调整为：

initMap <<<numBlocks, numThr>>> (map, pitch, width, height, 1, 1);

请注意，任何“幻像”元素在行上的索引大于w 将不会被初始化。它们将包含垃圾。但是您不应该担心这一点或使用这些元素。

当你将此数组复制回主机时，你应该使用cudaMemcpy2D，如果设置正确，它将在复制回主机的过程中消除幻像元素，因此主机数组中应该没有垃圾数据. SO上有很多问题展示了如何正确使用cudaMallocPitch/cudaMemcpy2D。

【讨论】：