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

Posted 2023-04-15

【中文标题】实验找出块大小对 cuda 程序速度的影响

【英文标题】：Experiment to find out affect of block size on cuda program speed

【发布时间】：2017-02-15 18:17:11

【问题描述】：

我想了解块中的线程数如何影响 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。” p>

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

【问题讨论】：

您只是定时启动开销。您实际上并没有计时内核的持续时间。在cudaPeekAtLastError 调用之后，添加对cudaDeviceSynchronize() 的调用，这将强制整个内核持续时间出现在您的计时中。

@RobertCrovella 是的！这就是问题所在。请写下您的答案，以便我将其标记为正确答案。

【参考方案1】：

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

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