CUDA向Kernel传递含矩阵的结构体时崩溃、返回全0问题排查

问题现象 #

• 代码无编译报错，但执行时异常崩溃、运行失控
• 所有计算返回值均为0，时钟计数溢出
• 调整代码后Kernel内嵌套else分支逻辑未按预期执行，返回值仍全为0

原问题代码 #

#define ROWS 700
#define COLS 1244
struct sobel {
    int Gradient[ROWS][COLS];
    int Image_input[ROWS][COLS];
    int G_x[ROWS][COLS];
    int G_y[ROWS][COLS];
};

__global__ void sobel(struct sobel* data)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    int XLENGTH = ROWS;
    int YLENGTH = COLS;

    if ((x < XLENGTH) && (y < YLENGTH))
    {
        if (x == 0 || x == XLENGTH - 1 || y == 0 || y == YLENGTH - 1)
        {
            data->G_x[x][y] = data->G_y[x][y] = data->Gradient[x][y] = 0;
        }
        else
        {
                data->G_x[x][y] = data->Image_input[x + 1][y - 1]
                    + 2 * data->Image_input[x + 1][y]
                    + data->Image_input[x + 1][y + 1]
                    - data->Image_input[x - 1][y - 1]
                    - 2 * data->Image_input[x - 1][y]
                    - data->Image_input[x - 1][y + 1];

                data->G_y[x][y] = data->Image_input[x - 1][y + 1]
                    + 2 * data->Image_input[x][y + 1]
                    + data->Image_input[x + 1][y + 1]
                    - data->Image_input[x - 1][y - 1]
                    - 2 * data->Image_input[x][y - 1]
                    - data->Image_input[x + 1][y - 1];

                data->Gradient[x][y] = abs(data->G_x[x][y]) + abs(data->G_y[x][y]);
                if (data->Gradient[x][y] > 255) {
                    data->Gradient[x][y] = 255;
                }
            
        }
    }
}

int main() {
    struct sobel* data = (struct sobel*)calloc(sizeof(*data), 1);
    struct sobel* dev_data; 
    cudaMalloc((void**)&dev_data, sizeof(*data));
    cudaMemcpy(dev_data, data, sizeof(data), cudaMemcpyHostToDevice);
    dim3 blocksize(16, 16);
    dim3 gridsize;
    gridsize.x = (ROWS + blocksize.x - 1) / blocksize.x;
    gridsize.y = (COLS + blocksize.y - 1) / blocksize.y;
    sobel <<< gridsize, blocksize >>> (dev_data);
    cudaMemcpy(data, dev_data, sizeof(data), cudaMemcpyDeviceToHost);
    free(data);
    cudaFree(dev_data);
    return 0;
}

问题解答 #

是否需要为结构体中每个矩阵单独分配设备内存？ #

不需要。
当前结构体内是完全连续的静态数组布局，cudaMalloc一次性分配整个结构体大小的设备内存是完全合法可用的。只有当结构体成员是独立指针（比如声明为int* Gradient、int* Image_input这类指针类型，指向分散的内存块）时，才需要为每个指针指向的内存区域单独做设备端分配，内嵌静态数组不需要单独分配。

核心错误点 #

1. cudaMemcpy拷贝长度完全错误
   两处cudaMemcpy调用传的长度参数都是sizeof(data)，data是主机端的结构体指针，64位系统下指针大小仅为8字节，而结构体总大小约为13.3MB（47001244*4字节）。仅往设备端拷贝8字节内容时，设备端结构体绝大多数内存都是未初始化的随机值，Kernel访问时直接触发越界崩溃；回拷时也只拷8字节回主机，自然读不到有效计算结果。
   修正方式：把两处cudaMemcpy的长度参数从sizeof(data)改为sizeof(*data)，确保完整拷贝整个结构体。

2. 输入数据未初始化
   用calloc分配主机端结构体内存时，所有字节会被初始化为0，之后没有给Image_input成员填充任何实际的图像像素值就直接拷贝到设备端。就算拷贝长度正确，输入全0的情况下Sobel算子计算出来的Gx、Gy、Gradient结果全为0，看起来就像else分支没有执行。
   修正方式：主机到设备拷贝前，先给data->Image_input填充实际待处理的像素数据。

3. 缺失CUDA错误检查
   代码没有对任何CUDA API调用、Kernel启动做返回值检查，内存拷贝越界、Kernel启动失败这类问题都会静默发生，很难快速定位。建议每次调用CUDA接口后都检查错误码，Kernel启动后加cudaDeviceSynchronize()再检查错误，避免运行失控找不到原因。

内容的提问来源于stack exchange，提问作者AlexandrosS