CUDA核函数For循环计算结果异常，请求排查原因

问题描述 #

编写了一段包含二维CUDA核函数的代码，预期执行后rr数组所有元素值应为33，但实际运行结果全为0。将Block配置改为(8,8,1)、Grid配置改为(row/8+1,col/8+1,1)后，部分元素结果正确，但仍有512个元素值为0。代码如下：

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

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#define row 65
#define col 13824

__global__ void tt(int *pp){
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    for(unsigned ig=0;ig<33;ig++){
        pp[i*col+j]+=1;
    }
    return;
}

int main(){

    int *pp;
    int *rr;

    pp=(int*)malloc(sizeof(int)*col*row);
    rr=(int*)malloc(sizeof(int)*col*row);

    memset(pp,0,sizeof(int)*row*col);

    int *pp_g;
    cudaMalloc((void**)&pp_g,sizeof(int)*row*col);
    cudaMemcpy(pp_g,pp,sizeof(int)*row*col,cudaMemcpyHostToDevice);

    dim3 block(32,32,1);
    dim3 grid(row/32+1,col/32+1,1);

    tt<<<grid,block>>>(pp_g);
    cudaDeviceSynchronize();

    cudaMemcpy(rr,pp_g,sizeof(int)*row*col,cudaMemcpyDeviceToHost);

    int ct=0;
    for(unsigned i=0;i<row*col;i++){
        if(rr[i]!=33){
            ct++;
        }
    }
    printf("%d\n",ct);

    return 0;

}

问题原因分析 #

1. 缺少线程边界检查
   核函数中没有判断线程索引i是否小于row、j是否小于col，导致超出有效范围的线程执行了数组写入操作：

   • 当i >= row或j >= col时，访问pp[i*col+j]要么越界到非法内存区域，触发CUDA运行时错误，导致核函数执行失败（表现为所有元素仍为初始值0）；
   • 若越界索引仍落在数组内存范围内（比如j=col时，i*col+j = (i+1)*col，对应第i+1行第0列的元素），会错误地修改其他有效元素的值，导致数据混乱。

2. Grid维度计算错误

   • 对于能被Block维度整除的维度（如col=13824，Block.y为32或8），Grid维度额外+1会产生多余的Block，导致大量线程的j超出col-1的有效范围；
   • 对于不能被Block维度整除的维度（如row=65），Grid维度+1是正确的向上取整方式，但必须配合边界检查过滤无效线程，否则会产生i >= row的无效线程。

3. 未添加CUDA错误检查
   代码中没有对cudaMalloc、cudaMemcpy、核函数启动、cudaDeviceSynchronize等操作进行错误检查，无法及时定位核函数执行失败的具体原因。

修复方案 #

修改后的代码 #

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

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#define row 65
#define col 13824

// CUDA错误检查宏
#define CHECK_CUDA_ERROR(err) \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA错误位于 %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    }

__global__ void tt(int *pp){
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    // 仅处理有效索引范围内的线程
    if (i < row && j < col) {
        for(unsigned ig=0;ig<33;ig++){
            pp[i*col+j] += 1;
        }
    }
}

int main(){
    int *pp = (int*)malloc(sizeof(int)*col*row);
    int *rr = (int*)malloc(sizeof(int)*col*row);
    memset(pp, 0, sizeof(int)*row*col);

    int *pp_g;
    cudaError_t err = cudaMalloc((void**)&pp_g, sizeof(int)*row*col);
    CHECK_CUDA_ERROR(err);

    err = cudaMemcpy(pp_g, pp, sizeof(int)*row*col, cudaMemcpyHostToDevice);
    CHECK_CUDA_ERROR(err);

    dim3 block(32, 32, 1);
    // 通用向上取整方式计算Grid维度，兼容整除与非整除场景
    dim3 grid((row + blockDim.x - 1)/blockDim.x, (col + blockDim.y - 1)/blockDim.y, 1);

    tt<<<grid, block>>>(pp_g);
    // 检查核函数启动错误
    err = cudaGetLastError();
    CHECK_CUDA_ERROR(err);

    err = cudaDeviceSynchronize();
    CHECK_CUDA_ERROR(err);

    err = cudaMemcpy(rr, pp_g, sizeof(int)*row*col, cudaMemcpyDeviceToHost);
    CHECK_CUDA_ERROR(err);

    int ct = 0;
    for(unsigned i=0; i<row*col; i++){
        if(rr[i] != 33){
            ct++;
        }
    }
    printf("错误元素数量: %d\n", ct);

    // 释放内存
    free(pp);
    free(rr);
    cudaFree(pp_g);

    return 0;
}

关键修改点 #

• 添加边界检查：核函数中通过if (i < row && j < col)确保只有有效线程执行数组修改操作；
• 修正Grid维度计算：使用(维度值 + Block维度 - 1)/Block维度的通用向上取整方式，避免多余Block导致的索引越界；
• 添加错误检查：通过CHECK_CUDA_ERROR宏检查所有CUDA操作的错误，便于调试；
• 内存释放：添加了主机端和设备端的内存释放操作，避免内存泄漏。

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