具有不同块大小的 CUDA 矩阵加法时序答案

【问题标题】：CUDA Matrix Addition Timing with varying block size具有不同块大小的 CUDA 矩阵加法时序
【发布时间】：2015-09-19 20:37:12
【问题描述】：

#include<stdio.h>
#include<cuda_runtime.h>
#include<sys/time.h>

#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);
   }
}

double measure_time()
{
struct timeval tp;
gettimeofday(&tp,NULL);
return ((double)tp.tv_sec+(double)tp.tv_usec*1.e-6);


}
__global__ void sum_matrix(int *a,int *b,int *c,int nx,int ny)
{

int ix=blockIdx.x*blockDim.x+threadIdx.x;
int iy=blockIdx.y*blockDim.y+threadIdx.y;
int idx=iy*nx+ix;
c[idx]=a[idx]+b[idx];
}

int main(int argc, char *argv[])
{
int dimx=atoi(argv[1]);
int dimy=atoi(argv[2]);

int nx=4096;
int ny=4096;

dim3 block (dimx,dimy);
dim3 grid (nx/dimx,ny/dimy);

double start,end;

int *a,*b,*c;
long long nbytes=nx*ny*sizeof(int);

cudaMalloc((int**)&a,nbytes);
cudaMalloc((int**)&b,nbytes);
cudaMalloc((int**)&c,nbytes);

start=measure_time();
sum_matrix<<<grid,block>>>(a,b,c,nx,ny);
cudaDeviceSynchronize();
gpuErrchk( cudaPeekAtLastError() );
end=measure_time();

printf("Time elapsed = %f ms\n",(end-start)*1000);

cudaFree(a);
cudaFree(b);
cudaFree(c);
return 0;
}

上面是一个二维矩阵加法内核，我用它来检查 MSI GTX 750 1 GB GDDR5 卡上不同块大小配置的执行时间。以下是不同块大小配置的执行时间结果。

./sum_matrix 32 32 经过的时间 = 3.028154 毫秒

./sum_matrix 32 16 经过的时间 = 3.180981 毫秒

./sum_matrix 16 32 经过的时间 = 2.942085 毫秒

./sum_matrix 16 16 经过的时间 = 3.238201 毫秒

./sum_matrix 64 8 经过的时间 = 3.020048 毫秒

./sum_matrix 64 16 经过的时间 = 3.304005 毫秒

./sum_matrix 128 2 经过的时间 = 2.965927 毫秒

./sum_matrix 128 1 经过的时间 = 2.896070 毫秒

./sum_matrix 256 2 经过的时间 = 3.004074 毫秒

./sum_matrix 256 1 经过的时间 = 2.948046 毫秒

我可以理解的是，将块大小增加到最大（1024 个线程），例如 (64,16) 可能会降低可用并行度，因此性能会更差。我不明白为什么增加块 x 尺寸和减少块 y 会提供更好的性能。是由于内存合并/缓存还是发散？

谢谢

【问题讨论】：

您在动态时钟游戏卡上运行。这使得计时结果非常不稳定。

标签： cuda gpu

【解决方案1】：

我认为您的主要问题是，这些差异一开始在统计上并不显着。对于如此少量的数据，实际执行内核启动的开销很可能占据了执行时间。请注意，无论使用的块大小如何，您的所有时间都在 3 毫秒左右。

您可以通过在循环中多次启动内核并平均执行时间来获得更精确的结果，但是对于如此小的内核调用，这可能只能用于确认所有启动都在大约由于启动和块调度开销支配了实际内核执行时间，所以时间量相同。

为了看到使用不同块大小的任何具有统计意义的结果，您可能需要做一些比仅仅 1600 万个整数加法更重要的事情。

【讨论】：

我怀疑内核启动开销是一个因素 - 它通常是 10us 的数量级。

【解决方案2】：

首先，正如我在评论中提到的，在游戏卡上计时时，尤其是对于如此短的测试，结果会出现波动。它们是动态时钟的，并且时钟不会在每次运行中都相同。

根据我的经验，块大小/形状不太可能对这种天真的实现的元素问题产生太大影响。只要您的块 x 尺寸是 32 的倍数，并且您有足够大的块来获得 100% 的占用率。在那之后，它只是流入和流出数据。

但是，您可以做得比您的实施更好。 This reference 现在已经很老了，但也有一些好处。从本质上讲，计算每个线程的许多元素将为您提供更好的性能。

另一个（小的）性能改进将来自对内存事务的矢量化。如果您对 int4s 而不是 int 进行操作，则硬件可以发出 128 字节的加载/存储指令，而不是 32 字节的加载/存储指令。这会稍微高效一些，因为要处理的内存指令少了 4 倍。

【讨论】：