CUDA向量加法：扩展线程网格后效率未提升的原因探究

我在测试CUDA加速向量加法时，从add<<<1, 1>>>(N, x, y);调整为add<<<1, 256>>>(N, x, y);，效率大幅提升——通过nvprof测试，耗时从65.409ms降至3.0164ms。

对应的代码如下：

#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  add<<<1, 256>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  cudaFree(x);
  cudaFree(y);
  
  return 0;
}

但当我使用包含多线程块的完整线程网格时，效率并未得到提升。

对应的多线程块版本代码：

#include <iostream>
#include <math.h> // function to add the elements of two arrays


__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
} 

int main(void)
{ 
  int N = 1<<20; 
  float *x, *y;
  cudaMallocManaged(&x, N * sizeof(float));
  cudaMallocManaged(&y, N * sizeof(float)); 
  
  
  // initialize x and y arrays on the host 
  for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } 


  // Run kernel on 1M elements on the GPU
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();
  
  // Check for errors (all values should be 3.0f) 
  float maxError = 0.0f; 
  for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i]-3.0f)); 
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory 
  cudaFree(x);
  cudaFree(y);

  return 0;
}

环境信息 #

• 设备：NVIDIA GeForce RTX 2080 Ti
• CUDA版本：12.6
• NVIDIA-SMI：560.35.03
• 驱动版本：560.35.03

按照CUDA官方教程内容，使用grid-stride loop应该能大幅提升运算速度，但实际测试结果不符合预期。

请问为何扩展线程网格后，CUDA向量加法的效率未得到提升？

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