CUDA C矩阵加法代码运行正常，nvprof无法检测内核函数如何解决？

1. 修复代码逻辑错误（核心原因） #

你当前代码的致命问题是直接向CUDA内核传入主机指针调用：

• 代码中matrixInit<<<>>>调用时传入的A/B/C是用malloc分配的主机内存指针，CUDA内核无法直接访问主机地址，这3次内核调用会直接启动失败
• 你没有添加任何CUDA错误检查逻辑，因此无法感知到内核启动失败的问题，误以为程序正常运行，实际上所有内核都没有成功执行，nvprof自然检测不到内核

修复方案 #

调整初始化逻辑，直接用内核初始化设备内存即可，同时添加错误检查确认内核运行状态：

#include <stdio.h>
#include <cuda.h>
#include <time.h>
#include <cuda_profiler_api.h>

__global__ void matrixInit(float *m, int N_1, int N_2, int value){
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int strideX = blockDim.x * gridDim.x;
    unsigned int strideY = blockDim.y * gridDim.y;

    for(int j=iy; j<N_2; j+=strideY){
        for(int i=ix; i<N_1; i+=strideX){
            m[j*N_1+i] = value;
        }
    }
}

__global__ void matrixAdd(float *d_A, float *d_B, float *d_C, int N_1, int N_2){
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int strideX = blockDim.x * gridDim.x;
    unsigned int strideY = blockDim.y * gridDim.y;

    for(int j=iy; j<N_2; j+=strideY){
        for(int i=ix; i<N_1; i+=strideX){
            d_C[j*N_1+i] = d_A[j*N_1+i]+d_B[j*N_1+i];
        }
    }
}

int main() {
    int N_1 = 1 << 12;
    int N_2 = 1 << 15;
    int N_1_2 = N_1 * N_2;
    size_t bytes = N_1_2*sizeof(float);

    float *A, *B, *C;
    float *d_A, *d_B, *d_C;
    A = (float*)malloc(bytes);
    B = (float*)malloc(bytes);
    C = (float*)malloc(bytes);

    // 先分配设备内存
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    int threadsPerBlock=32;
    dim3 threads(threadsPerBlock,threadsPerBlock);
    dim3 numBlocks( N_1/threads.x, N_2/threads.y);
    printf(" Grid Size of X: %d Grid Size of Y: %d \n ",threads.x,threads.y);

    // 直接用内核初始化设备内存
    matrixInit<<<numBlocks,threads>>>(d_A,N_1, N_2, 1.0f);
    matrixInit<<<numBlocks,threads>>>(d_B,N_1, N_2, 2.0f);
    matrixInit<<<numBlocks,threads>>>(d_C,N_1, N_2, 0.0f);
    // 检查内核启动错误
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("内核启动失败: %s\n", cudaGetErrorString(err));
        return 1;
    }
    // 同步等待内核执行完成
    cudaDeviceSynchronize();

    clock_t t = clock();
    matrixAdd<<<numBlocks,threads>>>(d_A, d_B, d_C, N_1, N_2);
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("matrixAdd启动失败: %s\n", cudaGetErrorString(err));
        return 1;
    }
    cudaDeviceSynchronize();

    cudaMemcpy(C, d_C, bytes, cudaMemcpyDeviceToHost);
    t = clock() - t;
    printf("Program executed at %f seconds", ((float)t) / CLOCKS_PER_SEC);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(A);
    free(B);
    free(C);

    cudaProfilerStop();
    return 0;
}

额外提示：你原本的matrixAdd内核存在逻辑错误，输出地址写为d_C[i]会导致不同j循环的结果互相覆盖，上述代码已经修正为d_C[j*N_1+i]。

2. 调整nvprof运行参数解决缓冲区不足警告 #

按照警告提示调大设备缓冲区和信号量池大小，运行命令改为：

nvprof --device-buffer-size 128MB --profiling-semaphore-pool-size 16384 ./你的可执行文件名

3. 高版本CUDA/新架构GPU适配 #

如果你使用的是安培及以上架构GPU（GeForce 30系、40系等）+ CUDA 11及以上版本，nvprof已经不再支持内核性能分析，官方替换工具为Nsight Systems，使用以下命令即可完成性能分析：

nsys profile ./你的可执行文件名

执行完成后会生成.qdrep格式的报告文件，可用Nsight Systems GUI工具打开查看内核性能数据。

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