实验找出块大小对 cuda 程序速度的影响

Question

我想了解块中的线程数如何影响 cuda 程序的性能和速度。我写了一个简单的矢量加法代码，这是我的代码：

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}
__global__ void gpuVecAdd(float *a, float *b, float *c, int n) {
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    if (id < n) {
        c[id] = a[id] + b[id];
    }
}
int main() {
    int n = 1000000;
    float  *h_a, *h_b, *h_c, *t;
    srand(time(NULL));
    size_t bytes = n* sizeof(float);
    h_a = (float*) malloc(bytes);
    h_b = (float*) malloc(bytes);
    h_c = (float*) malloc(bytes);
    for (int i=0; i<n; i++)
        {
         h_a[i] =rand()%10;
          h_b[i] =rand()%10;
        }
   float *d_a, *d_b, *d_c;
   cudaMalloc(&d_a, bytes);
   cudaMalloc(&d_b, bytes);
   cudaMalloc(&d_c, bytes);

   gpuErrchk( cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
   gpuErrchk( cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

   clock_t t1,t2;
   t1 = clock();
   int block_size = 1024;
   gpuVecAdd<<<ceil(float(n/block_size)),block_size>>>(d_a, d_b, d_c, n);
   gpuErrchk( cudaPeekAtLastError() );
   t2 = clock();
   cout<<(float)(t2-t1)/CLOCKS_PER_SEC<<" seconds";
   gpuErrchk(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
cudaFree(d_a);
   cudaFree(d_b);
   cudaFree(d_c);
   free(h_a);
   free(h_b);
   free(h_c);
}

我阅读了 this post 并根据 talonmies 的回答“每个块的线程数应该是 warp 大小的整数倍，在所有当前硬件上都是 32。"

我用每个块的不同线程数检查了代码，例如 2 和 1024（这是 32 的乘积，也是每个块的最大线程数）。两种尺寸的平均运行时间几乎相等，我看不出它们之间存在巨大差异。这是为什么？我的基准测试不正确吗？

Answer 1

CUDA 中的 GPU 内核启动是异步。这意味着在内核完成执行之前，控制权将返回到 CPU 线程。

如果我们希望 CPU 线程对内核的持续时间进行计时，我们必须让 CPU 线程等待内核完成。我们可以通过在计时区域调用 cudaDeviceSynchronize() 来做到这一点。然后测量的时间将包括内核执行的完整持续时间。

实验找出块大小对 cuda 程序速度的影响

Experiment to find out affect of block size on cuda program speed

cuda

gpu

gpgpu

nvidia