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

Question

#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 毫秒

我能理解的是，像 (64,16) 那样将块大小增加到最大值（1024 个线程）可能会减少可用的并行度，因此性能会更差。我不明白为什么增加块 x 维度和减少块 y 可以提供更好的性能。是记忆力coalescing/caching还是分歧？

谢谢

Answer 1

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

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

为了查看使用不同块大小的任何具有统计显着性的结果，您可能需要做一些（远）比仅仅 1600 万整数加法更重要的事情。

Answer 2

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

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

但是，您可以比您的实施做得更好。 This reference 现在已经很老了，但提出了一些很好的观点。本质上，每个线程计算很多元素会给你更好的性能。

另一个（小的）性能改进将来自矢量化内存事务。如果你对 int4s 而不是 ints 进行操作，硬件可以发出 128 字节 load/store 指令而不是 32 字节 load/store 指令。由于要处理的内存指令减少了 4 倍，因此效率稍高。

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

CUDA Matrix Addition Timing with varying block size

cuda

gpu