CUDA Kernel有效带宽计算异常：结果远超理论值求修正

问题背景 #

尝试计算以下CUDA Kernel的有效带宽时，得到了远超设备理论带宽（448GB/s）的1.5TB/s结果：

__global__ void compute_flags(int nquarts, int4 *d_flag, int4 * dat_arr, int * arr, int array_size){
    int gi = blockIdx.x * blockDim.x + threadIdx.x;
    if (gi < nquarts) {
        int4 res = {0, 0, 0, 0};
        int4 dat = dat_arr[gi];

        for (int i = 0; i < array_size; i++){
            int p = arr[i];
            res.x = res.x | (p == dat.x);
            res.y = res.y | (p == dat.y);
            res.z = res.z | (p == dat.z);
            res.w = res.w | (p == dat.w);
        }
        d_flag[gi] = res;
    }
}

当前使用的字节访问量计算公式为：

2 * n_kernel_calls * sizeof(int4) + n_kernel_calls * array_size * sizeof(int)

完整测试代码如下：

#include <iostream>
#include <stdlib.h>

#define arr_size 1000000

int round_div_up (int a, int b){
    return (a + b - 1)/b;
}

void cuda_err_check (cudaError_t err, const char *file, int line)
{
    if (err != cudaSuccess)
    {
        fprintf (stderr, "CUDA error: %s (%s:%d)\n", cudaGetErrorString (err), file, line);
        exit (EXIT_FAILURE);
    }
}

__global__ void compute_flags(int nquarts, int4 *d_flag, int4 * dat_arr, int * arr, int array_size){
    int gi = blockIdx.x * blockDim.x + threadIdx.x;
    if (gi < nquarts) {
        int4 res = {0, 0, 0, 0};
        int4 dat = dat_arr[gi];

        for (int i = 0; i < array_size; i++){
            int p = arr[i];
            res.x = res.x | (p == dat.x);
            res.y = res.y | (p == dat.y);
            res.z = res.z | (p == dat.z);
            res.w = res.w | (p == dat.w);
        }
        d_flag[gi] = res;
    }
}

using namespace std;

int main(void){
    int V1 [arr_size] = {};
    int V2 [arr_size] = {};

    // fill with random numbers
    for(int i = 0; i < arr_size; i++){
        V1[i] = rand() % 100;
        V2[i] = rand() % 100;
    }
    
    int4 *d_flag;
    int * d_V1;
    int * d_V2;
    cudaError_t err;
    cudaEvent_t start, stop;

    err = cudaEventCreate(&start); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaEventCreate(&stop); cuda_err_check(err, __FILE__, __LINE__);

    err = cudaMalloc((void **)&d_flag, arr_size * sizeof(int)); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaMalloc((void **)&d_V1, arr_size * sizeof(int)); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaMalloc((void **)&d_V2, arr_size * sizeof(int)); cuda_err_check(err, __FILE__, __LINE__);

    err = cudaMemcpy(d_V1, V1, arr_size * sizeof(int), cudaMemcpyHostToDevice); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaMemcpy(d_V2, V2, arr_size * sizeof(int), cudaMemcpyHostToDevice); cuda_err_check(err, __FILE__, __LINE__);

    uint64_t nquarts = round_div_up(arr_size, 4);
    uint64_t lws = 256;
    uint64_t gws = round_div_up(nquarts, lws);

    err = cudaEventRecord(start); cuda_err_check(err, __FILE__, __LINE__);

    compute_flags<<<gws, lws>>>(nquarts, d_flag, (int4*)d_V1, d_V2, arr_size);

    err = cudaEventRecord(stop); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaEventSynchronize(stop); cuda_err_check(err, __FILE__, __LINE__);

    err = cudaGetLastError(); cuda_err_check(err, __FILE__, __LINE__);
    err = cudaDeviceSynchronize(); cuda_err_check(err, __FILE__, __LINE__);

    uint64_t byte_accesses = 2 * nquarts * sizeof(int4) + nquarts * arr_size * sizeof(int);
    float time;

    err = cudaEventElapsedTime(&time, start, stop); cuda_err_check(err, __FILE__, __LINE__);

    cout << "Time: " << time << " ms" << endl;
    cout << "Bandwidth: " << byte_accesses / time / 1e6 << " GB/s" << endl;

}

核心错误分析 #

1. 全局内存访问量重复计算
   公式中n_quarts * array_size * sizeof(int)的逻辑错误：循环中所有线程读取的是同一个全局数组arr，CUDA的L1/L2缓存会缓存该数组数据，全局内存只需要读取一次arr的全部内容，而非每个线程都完整读取一遍。这是导致计算出的带宽远超理论值的根本原因。

2. 内存分配类型不匹配（代码bug）
   主函数中d_flag是int4*类型，但分配内存时使用了arr_size * sizeof(int)，实际需要的是nquarts * sizeof(int4)，虽然这不会直接影响带宽计算，但会导致内存越界风险。

正确的带宽计算方法 #

有效带宽的计算基于实际从全局内存读取/写入的字节总量，正确的字节访问量公式为：

uint64_t byte_accesses = nquarts * sizeof(int4)  // 读取dat_arr的字节数
                       + nquarts * sizeof(int4)  // 写入d_flag的字节数
                       + arr_size * sizeof(int); // 读取arr的字节数（仅一次全局读取）

简化后为：

uint64_t byte_accesses = 2 * nquarts * sizeof(int4) + arr_size * sizeof(int);

公式验证示例 #

以arr_size=1e6为例，nquarts=250000（1e6/4），sizeof(int4)=16，sizeof(int)=4：
总字节数 = 225000016 + 1e6*4 = 8,000,000 + 4,000,000 = 12,000,000字节（12MB）
假设kernel运行时间为0.027ms，带宽计算为：
12e6 / (0.027e-3) / 1e9 ≈ 444 GB/s，与设备理论带宽448GB/s接近，结果合理。

额外代码修正 #

修正d_flag的内存分配语句：

err = cudaMalloc((void **)&d_flag, nquarts * sizeof(int4)); cuda_err_check(err, __FILE__, __LINE__);

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