对于动态内存分配的二维数组，使用cudaMallocPitch和cudaMemcpy2D进行内存分配与拷贝后，程序运行崩溃

Question

下面是我的问题代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"

#include <stdio.h>
#include <stdlib.h>


#define W 25
#define H 25

// 核函数
__global__ void kernel(int* a, size_t pitch)
{
    int x = threadIdx.x;
    int y = threadIdx.y;
    int *row_a = (int*)((char*)a + y * pitch);
    // Clear to zero
    row_a[x] = 0;
}
int main()
{
    int **a;   
    int *dev_a;
    size_t pitch;
    dim3 threads(W, H);
    // 为a动态分配内存
    a = (int**)malloc(H * sizeof(int*));
    for (int i = 0; i < H; i++)
    {
        a[i] = (int*)malloc(W * sizeof(int));
    }
    // 初始化数组a
    for (int i = 0; i < H; i++)
    {
        for (int j = 0; j < W; j++)
        {
            a[i][j] = 1;
        }
    }
    printf("修改前的数组内容：\n");
    for (int i = 0; i < H; i++)
    {
        for (int j = 0; j < W; j++)
        {
            printf("%d ", a[i][j]);
        }
        printf("\n");
    }
    // 使用cudaMallocPitch分配设备内存
    cudaMallocPitch((void**)&dev_a, &pitch, W * sizeof(int), H);
    // 将数组a中的内容拷贝到设备数组dev_a上    
    cudaMemcpy2D(dev_a, pitch, a, W * sizeof(int), W * sizeof(int), H, cudaMemcpyHostToDevice);
    // 调用核函数
    kernel<<<1, threads>>>(dev_a, pitch);
    // 将结果拷贝回主机
    cudaMemcpy2D(a, W * sizeof(int), dev_a, pitch, W * sizeof(int), H, cudaMemcpyDeviceToHost);
    printf("修改后的数组内容：\n");
    for (int i = 0; i < H; i++)
    {
        for (int j = 0; j < W; j++)
        {
            printf("%d ", a[i][j]);
        }
        printf("\n");
    }

    return 0;
}

上面代码编译是没有错误的，但是运行时程序会崩溃。当我将动态数组改成静态数组之后，就不会出现这个问题了。希望各位能帮忙解决一下。

Answer 1

我在stackoverflow找到了类似的问题，原来是因为cudaMemcpy2D函数拷贝的内存必须是连续的。如果采用一般的动态内存分配的方法，就像我的代码中那样，

// 为a动态分配内存
    a = (int**)malloc(H * sizeof(int*));
    for (int i = 0; i < H; i++)
    {
        a[i] = (int*)malloc(W * sizeof(int));
    }

这样的话，我们得到的二维数组在内存中并不是连续的，所以拷贝到设备上就是一段错误的内存空间，运行时就会出错。当然动态的二维数组也可以有其他方法使其连续,即先分配一段连续的内存空间，然后再创建一个指针数组，最后将这些指针按照指定的步长(宽度)指向对应的内存地址即可：

float* bigArray = new float[h * w]
float** myArray2 = new float*[h]
for( int i = 0 ; i < h ; i++ )
{
   myArray2[i] = &bigArray[i * w];
}

当然，我们也可以直接使用上面这一段连续的内存空间(bigArray)作为我们的“二维数组”，只要我们以步长(宽度)的方式访问数组就可以了：

for (int i = 0; i < h; ++i)
{
    // 在这里w宽度即为步长(stride)
    for (int j = 0; j < w; ++j)
        printf("%lf ", bigArray[i * w + j]);
}