c++ - CUDA 周期性执行时间

Question

我刚开始学习 CUDA，在解释我的实验结果时遇到了麻烦。我想在一个将两个 vector 相加的简单程序中比较 CPU 与 GPU。代码如下:

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

void add_cpu(int* a, int* b, int* c, long long n) {
    for (long long i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

void check_results(int* gpu, int* cpu, long long n) {
    for (long long i = 0; i < n; i++) {
        if (gpu[i] != cpu[i]) {
            printf("Different results!\n");
            return;
        }
    }
}

int main(int argc, char* argv[]) {
    long long n = atoll(argv[1]);
    int num_of_blocks = atoi(argv[2]);
    int num_of_threads = atoi(argv[3]);

    int* a = new int[n];
    int* b = new int[n]; 
    int* c = new int[n]; 
    int* c_cpu = new int[n];
    int *dev_a, *dev_b, *dev_c;
    cudaMalloc((void **) &dev_a, n * sizeof(int));
    cudaMalloc((void **) &dev_b, n * sizeof(int));
    cudaMalloc((void **) &dev_c, n * sizeof(int));
    for (long long i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }

    cudaMemcpy(dev_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, n * sizeof(int), cudaMemcpyHostToDevice);

    StopWatchInterface *timer=NULL;
    sdkCreateTimer(&timer);
    sdkResetTimer(&timer);
    sdkStartTimer(&timer);  

    add <<<num_of_blocks, num_of_threads>>>(dev_a, dev_b, dev_c, n);

    cudaDeviceSynchronize();
    sdkStopTimer(&timer);
    float time = sdkGetTimerValue(&timer);
    sdkDeleteTimer(&timer);

    cudaMemcpy(c, dev_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    clock_t start = clock();
    add_cpu(a, b, c_cpu, n);
    clock_t end = clock();

    check_results(c, c_cpu, n);
    printf("%f %f\n", (double)(end - start) * 1000 / CLOCKS_PER_SEC, time);

    return 0;
}

我使用 bash 脚本循环运行此代码:

for i in {1..2560}
do
    n="$((1024 * i))"
    out=`./vectors $n $i 1024`
    echo "$i $out" >> "./vectors.txt"
done

其中 2560 是我的 GPU 支持的最大块数，1024 是块中的最大线程数。所以我只是将它的最大块大小运行到我的 GPU 可以处理的最大问题大小，步长为 1 个块( vector 中为 1024 个整数)。
这是我的 GPU 信息:

CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "NVIDIA GeForce RTX 2070 SUPER"
  CUDA Driver Version / Runtime Version          11.3 / 11.0
  CUDA Capability Major/Minor version number:    7.5
  Total amount of global memory:                 8192 MBytes (8589934592 bytes)
  (040) Multiprocessors, (064) CUDA Cores/MP:    2560 CUDA Cores
  GPU Max Clock rate:                            1785 MHz (1.78 GHz)
  Memory Clock rate:                             7001 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 4194304 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        65536 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  1024
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 11.3, CUDA Runtime Version = 11.0, NumDevs = 1
Result = PASS

运行实验后，我收集了结果并绘制了它们:

所以困扰我的是 GPU 执行时间中的 256 个块宽的时间段。我不知道为什么会发生这种情况。为什么执行 512 块比执行 513 块线程慢很多？
我还用恒定数量的块 (2560) 以及不同的块大小检查了这一点，它总是给出 256 * 1024 vector 大小的时间段(因此对于块大小 512，它每个 512 个块，而不是每个 256 个块)。所以也许这与内存有关，但我不知道是什么。
我将不胜感激关于为什么会发生这种情况的任何想法。

Answer 1

这绝不是一个完整或准确的答案。但是，我相信您观察到的周期性模式至少部分是由于 1 次或首次内核启动开销。良好的基准测试实践通常是做一些不同于你正在做的事情。例如，多次运行内核并取平均值。或者做一些其他类型的统计测量。
当我在 GTX 960 GPU 上使用您的脚本运行您的代码时，我得到以下图表(仅绘制 GPU 数据，垂直轴以毫秒为单位):

当我按如下方式修改您的代码时:

cudaMemcpy(dev_c, c, n * sizeof(int), cudaMemcpyHostToDevice);
// next two lines added:
add <<<num_of_blocks, num_of_threads>>>(dev_a, dev_b, dev_c, n);
cudaDeviceSynchronize();

StopWatchInterface *timer=NULL;
sdkCreateTimer(&timer);
sdkResetTimer(&timer);
sdkStartTimer(&timer);  

add <<<num_of_blocks, num_of_threads>>>(dev_a, dev_b, dev_c, n);

cudaDeviceSynchronize();

首先进行“热身”运行，然后为第二次运行计时，我见证了这样的数据:

所以没有预热的数据显示出周期性。热身后，周期性消失。我得出的结论是，周期性是由于某种 1-time 或 first-time 行为。可能属于这一类别的一些典型事物是缓存效果和 cuda“懒惰”初始化效果(例如，对 GPU 代码进行 JIT 编译所花费的时间，这在您的情况下肯定会发生，或者加载 GPU 代码的时间到 GPU 内存中)。我将无法进一步解释究竟是哪种首次效应导致了周期性。
另一个观察结果是，虽然我的数据显示了每个图形的预期“平均斜率”，表明与 2560 个块相关的内核持续时间大约是与 512 个块相关的内核持续时间的 5 倍，但我没有看到您的这种趋势数据。然而，它应该在那里。您的 GPU 将在大约 40 个块时“饱和”。此后，平均内核持续时间应该以近似线性的方式增加，使得与 2560 个块相关的内核持续时间是与 512 个块相关的内核持续时间的 4-5 倍。我根本无法在这方面解释您的数据，我怀疑是图形或数据处理错误，或者您的环境中的某个特征(例如与其他用户共享 GPU、损坏的 CUDA 安装等)在我的环境，我无法猜测。
最后，我的结论是，在良好的基准测试技术存在的情况下，GPU 的“预期”行为更加明显。