如何正确调试CUDA内核？解决Warp非法地址异常问题

问题描述 #

编写的CUDA代码调用broadcastMatrix时触发received signal CUDA_EXCEPTION_14, Warp Illegal Address错误。涉及主机到设备内存拷贝、自定义GPU内存分配、广播矩阵内核等逻辑。使用VSCode + launch.json + cuda-gdb调试时，可设置断点但无法单步调试内核，错误随机出现在第一个block的不同线程。此前其他内核运行正常，所有cudaMalloc均返回cudaSuccess，需解决异常并配置好内核调试环境，避免依赖printf调试。

代码片段 #

主机到设备内存拷贝函数 #

void initializeCudaMatrixWithCPUMatrix(float *hostA, float *cudaA, int n,
                                       int m) {
  // float testt = hostA[888];
  cudaMalloc((void **)&cudaA, n * m * sizeof(float));
  cudaMemcpy(cudaA, hostA, n * m * sizeof(float), 
  cudaMemcpyHostToDevice);
  cudaDeviceSynchronize();
  // testt = cudaA[888];
}

调用方式：

initializeCudaMatrixWithCPUMatrix(model->embeddingMatrix,
                                    model->embeddingMatrixCuda, 408, 1024);

自定义GPU内存分配函数 #

void gpuMallocMatrix(float **A, int m, int n) {
  cudaError_t cudaStatus = cudaMalloc((void **)A, n * m * sizeof(float));
}

broadcastedMatrix内存分配调用：

float *broadcastedEmbeddingMatrix;
gpuMallocMatrix(&broadcastedEmbeddingMatrix, trademodel->maxnum * 128, 1024);
cudaDeviceSynchronize();

广播矩阵函数及内核 #

extern "C" void broadcastMatrix( float *Matrix, float *BroadcastedMatrix,
                                 int rows, int cols, int batchsize )   {   
    dim3 rows2d(rows, batchsize);
    broadcastMatrixKernel<<<rows2d, cols>>>(Matrix, BroadcastedMatrix,
                                            rows, cols, batchsize); 
}

__global__ void broadcastMatrixKernel(float *Matrix, float *BroadcastedMatrix, 
                                      int rows, int cols, int batchsize) {
   int rowIdx = blockIdx.x;
   int batchIdx = blockIdx.y;
   int colIdx = threadIdx.x;
   if (rowIdx < rows && colIdx < cols && batchIdx < batchsize) { 
       BroadcastedMatrix[batchIdx * rows * cols + rowIdx * cols + colIdx] = 
       Matrix[rowIdx * cols + colIdx];   
   }
}

触发错误的调用：

broadcastMatrix(trademodel->embeddingMatrixCuda, 
                broadcastedEmbeddingMatrix, 408, 1024, 128); 

问题解决 #

1. 内存访问错误修复 #

（1）修复设备指针传递错误

initializeCudaMatrixWithCPUMatrix中cudaA按值传递，cudaMalloc仅修改局部变量，外部的model->embeddingMatrixCuda未被正确初始化，导致内核访问非法地址。修改为指针的指针传递：

void initializeCudaMatrixWithCPUMatrix(float *hostA, float **cudaA, int n,
                                       int m) {
  cudaError_t status = cudaMalloc((void **)cudaA, n * m * sizeof(float));
  if (status != cudaSuccess) {
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(status));
    exit(EXIT_FAILURE);
  }
  status = cudaMemcpy(*cudaA, hostA, n * m * sizeof(float), cudaMemcpyHostToDevice);
  if (status != cudaSuccess) {
    printf("cudaMemcpy failed: %s\n", cudaGetErrorString(status));
    exit(EXIT_FAILURE);
  }
  cudaDeviceSynchronize();
}

调用时传入指针地址：

initializeCudaMatrixWithCPUMatrix(model->embeddingMatrix,
                                    &model->embeddingMatrixCuda, 408, 1024);

（2）修正广播矩阵内存尺寸

当前gpuMallocMatrix调用中，内存尺寸与实际需求不匹配（需batchsize*rows*cols即1284081024个float），若trademodel->maxnum不等于408，会导致内存越界。修改分配调用：

gpuMallocMatrix(&broadcastedEmbeddingMatrix, 128*408, 1024);

同时给gpuMallocMatrix添加错误检查：

void gpuMallocMatrix(float **A, int m, int n) {
  cudaError_t cudaStatus = cudaMalloc((void **)A, n * m * sizeof(float));
  if (cudaStatus != cudaSuccess) {
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(cudaStatus));
    exit(EXIT_FAILURE);
  }
}

（3）添加内核启动错误检查

在broadcastMatrix中添加内核启动后的错误检测，提前发现配置问题：

extern "C" void broadcastMatrix( float *Matrix, float *BroadcastedMatrix,
                                 int rows, int cols, int batchsize )   {   
    dim3 rows2d(rows, batchsize);
    broadcastMatrixKernel<<<rows2d, cols>>>(Matrix, BroadcastedMatrix,
                                            rows, cols, batchsize); 
    cudaError_t status = cudaGetLastError();
    if (status != cudaSuccess) {
        printf("Kernel launch failed: %s\n", cudaGetErrorString(status));
        exit(EXIT_FAILURE);
    }
    cudaDeviceSynchronize();
}

2. VSCode + cuda-gdb内核调试配置 #

（1）编译选项配置

编译时必须添加-g -G参数生成调试信息，示例nvcc命令：

nvcc -g -G your_code.cu -o your_executable

若用CMake，在CMakeLists.txt中添加：

set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -g -G")

（2）launch.json配置

确保配置正确的程序路径与cuda-gdb路径，示例配置：

{
    "version": "0.2.0",
    "configurations": [
        {
            "name": "CUDA Debug",
            "type": "cppdbg",
            "request": "launch",
            "program": "${workspaceFolder}/your_executable",
            "args": [],
            "stopAtEntry": false,
            "cwd": "${workspaceFolder}",
            "environment": [],
            "externalConsole": false,
            "MIMode": "gdb",
            "miDebuggerPath": "/usr/bin/cuda-gdb", // 根据实际路径调整
            "setupCommands": [
                {
                    "description": "Enable pretty-printing for gdb",
                    "text": "-enable-pretty-printing",
                    "ignoreFailures": true
                }
            ]
        }
    ]
}

（3）调试注意事项

• 确保GPU计算能力≥3.2（支持内核调试）
• 调试时先缩小测试规模（如batchsize=2、rows=10），便于单步跟踪线程状态
• 在内核断点处，可通过VSCode调试面板查看blockIdx、threadIdx等变量，验证内存索引计算是否正确

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