CUDA向量加法核：手动计算与Nsight Compute Roofline结果不符

我实现了一个简单的CUDA向量加法核C = A + B，三个向量规模均为N。手动计算参数如下：

• 算术操作数：N次单精度浮点加法
• 内存访问字节数：3×N×4（假设sizeof(float)==4，包含2次全局加载、1次全局存储）
• 算术强度（浮点操作数/内存访问字节数）：约0.083
• GFLOP/s计算方式：N×1e-9 / 核函数运行时间(秒)

但Nsight Compute的Roofline分析显示算术强度为0.12，与手动计算值存在明显差异。已知Nsight Compute会降低设备时钟频率，但算术强度理论上不受频率影响，数值应基本一致。

查看Nsight Compute的指令统计，全局加载（LDG）与存储（STG）操作数总和是FADD操作数的3倍，与手动逻辑一致，但仍无法解释算术强度的差异。请问该差异的原因是什么？是我手动计算存在错误吗？

核函数代码 #

#include <iostream>
#include <cuda_runtime.h>

#define N 200000

__global__ void vectorAdd(float* a, float* b, float* c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N)
    {
        c[tid] = a[tid] + b[tid];
    }
}

int main()
{
    // Declare and initialize host vectors
    float* host_a = new float[N];
    float* host_b = new float[N];
    float* host_c = new float[N];
    for (int i = 0; i < N; ++i)
    {
        host_a[i] = i;
        host_b[i] = 2 * i;
    }

    // Declare and allocate device vectors
    float* dev_a, * dev_b, * dev_c;
    cudaMalloc((void**)&dev_a, N * sizeof(float));
    cudaMalloc((void**)&dev_b, N * sizeof(float));
    cudaMalloc((void**)&dev_c, N * sizeof(float));

    // Copy host vectors to device
    cudaMemcpy(dev_a, host_a, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, host_b, N * sizeof(float), cudaMemcpyHostToDevice);

    // Define kernel launch configuration
    int blockSize, gridSize;
    cudaOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, vectorAdd, 0, N);

    // Start timer
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start);

    // Launch kernel
    vectorAdd<<<gridSize, blockSize>>>(dev_a, dev_b, dev_c);

    // Stop timer and calculate execution duration
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    // Copy result from device to host
    cudaMemcpy(host_c, dev_c, N * sizeof(float), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    // Print execution duration
    std::cout << "Kernel execution duration: " << milliseconds << " ms" << std::endl;

    int numFloatingPointOps = N;
    int numBytesAccessed = 3 * N * sizeof(float);
    float opsPerByte = static_cast<float>(numFloatingPointOps) / static_cast<float>(numBytesAccessed);

    std::cout << "Floating-point operations per byte: " << opsPerByte << std::endl;

    float executionTimeSeconds = milliseconds / 1e3;
    float numGFLOPs = static_cast<float>(numFloatingPointOps) / 1e9;
    float GFLOPs = numGFLOPs / executionTimeSeconds;

    std::cout << "GFLOP/s: " << GFLOPs << std::endl;

    // Cleanup
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
    delete[] host_a;
    delete[] host_b;
    delete[] host_c;

    return 0;
}

本地运行示例输出 #

Kernel execution duration: 0.014144 ms
Floating-point operations per byte: 0.0833333
GFLOP/s: 14.1403

编译与Profiling命令 #

nvcc vectorAdd.cu
sudo env "PATH=$PATH" ncu -f -o vectorAdd_rep --set full ./a.out

问题解答 #

你的手动计算逻辑完全正确，差异的核心原因是Nsight Compute的Roofline分析对内存访问字节数的统计范围与你手动计算的不一致：

1. 统计范围差异
   你手动计算时包含了全局加载+全局存储的所有内存访问字节数（3×N×4），但Nsight Compute的Roofline分析默认可能仅统计全局加载的字节数，排除了存储操作的字节数。

   验证这个结论：如果仅计算加载字节数（2×N×4=1600000字节），算术强度为200000 / 1600000 = 0.125，与Nsight显示的0.12几乎一致，微小差异来自统计时的四舍五入或硬件执行的微小波动。

2. 为什么会有这种统计差异
   部分Roofline分析工具会优先关注加载带宽的瓶颈（因为加载通常是内存-bound任务的主要瓶颈），因此默认仅统计加载操作的内存访问。你可以在Nsight Compute的Roofline分析设置中，检查是否有选项可以调整内存访问的统计范围，确认是否包含存储操作。

3. 额外验证点
   你提到Nsight统计的LDG+STG操作数是FADD的3倍，说明工具确实统计了所有内存操作，但Roofline模块的算术强度计算逻辑单独排除了存储字节数，这是该工具的默认行为差异，而非你的计算错误。

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