CUDA理论带宽vs有效带宽

我有一个CUDA内核，它将两个矩阵相乘，其中Width和Height是我使用的块大小的倍数。CUDA理论带宽vs有效带宽

基于NVIDIA Quadro Fx的3800我使用了50 Gb/s的理论带宽，我有一些奇怪的结果（有效带宽比理论带宽大）

我将张贴在这里的一些结果：

随着块大小2

[10] [10] * [10] [10] - > BW = 0,02 Gb/s的[1000] [1000] * [1000] [1000] - > BW = 69,4 Gb/s

随着块大小64

[1000] [1000] * [1000] [1000] - > BW = 486,4 Gb/s的 [10000] [10000] * [10000] [10000] - > BW = 45072,12 Gb/s的

我从Nvidia的最佳实践指南的有效带宽的式（I已经简化，但它的等价物（除非有低级错误））。我认为内核很好，因为它与我阅读的一些Nvidia讲座非常相似（如果不相等），也因为它的正常工作（afaik）。

#define blocksize 64 
#define HM (10000) 
#define WM (10000) 
#define WN (10000) 
#define HN WM 
#define WP WN 
#define HP HM 
#define PTH WM 
#define PTW HM 

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN) 
    { 
__shared__ float MS[blocksize][blocksize]; 
__shared__ float NS[blocksize][blocksize]; 

int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y; 
int rowM=ty+by*blocksize; 
int colN=tx+bx*blocksize; 
int Pvalue=0; 

for(int m=0; m< uWM/blocksize;m++){ 
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)]; 
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)]; 
    __syncthreads(); 
    for(int k=0;k<blocksize;k++) 
     Pvalue+=MS[ty][k]*NS[k][tx]; 
    P[rowM*WP+colN]=Pvalue; 
} 

} 
int main(){ 


cudaEvent_t evstart, evstop; 
cudaEventCreate(&evstart); 
cudaEventCreate(&evstop); 

float*M=(float*)malloc(sizeof(float)*HM*WM); 
float*N=(float*)malloc(sizeof(float)*HN*WN); 

for(int i=0;i<WM*HM;i++) 
    M[i]=(float)i; 
for(int i=0;i<WN*HN;i++) 
    N[i]=(float)i; 




float*P=(float*)malloc(sizeof(float)*HP*WP); 



float *Md,*Nd,*Pd; 
cudaMalloc((void**)&Md,HM*WM*sizeof(float)); 

cudaMalloc((void**)&Nd,HN*WN*sizeof(float)); 

cudaMalloc((void**)&Pd,HP*WP*sizeof(float)); 



cudaMemcpy(Md,M,HM*WM*sizeof(float),cudaMemcpyHostToDevice); 

cudaMemcpy(Nd,N,HN*WN*sizeof(float),cudaMemcpyHostToDevice); 



dim3 dimBlock(blocksize,blocksize);//(tile_width , tile_width); 
dim3 dimGrid(WN/dimBlock.x,HM/dimBlock.y);//(width/tile_width , width/tile_witdh); 

cudaEventRecord(evstart,0); 

nonsquare<<<dimGrid,dimBlock>>>(Md,Nd,Pd,WM,WN); 

cudaEventRecord(evstop,0); 
cudaEventSynchronize(evstop); 
float time; 
cudaEventElapsedTime(&time,evstart,evstop); 

cudaMemcpy(P,Pd,WP*HP*sizeof(float),cudaMemcpyDeviceToHost); 

    cudaFree(Md); 
cudaFree(Nd); 
cudaFree(Pd); 


    printf("\ntime spent:%f",time); 
float Bandwidth=(HM*WM*4+WN*HN*4+HP*WP*4)/(time*1000000);/
printf("\nEffective Bandwidth:%f Gb/s\n",Bandwidth); 
    }

在此先感谢

来源

2011-03-22 Bernardo

你的问题是？ – 2011-03-22 18:23:21

有效带宽如何超过理论带宽？我认为理论带宽是显卡能达到的最大值，还是我错了？ – Bernardo 2011-03-22 18:54:11

我认为内核只是默默地失败。

您是否在内核调用后检查了任何错误？
代码是否工作？
你对时间有什么结果？

来源

2011-03-28 06:17:33 fabrizioM

我发现了这个问题。我没有关注每块的最大线程数。由于我使用64块大小，我不在乎，问题是它的2D块，所以64 * 64 = 4096线程每块。这是错误 – Bernardo 2011-03-28 14:18:06

我想到内核确实失败了:) – fabrizioM 2011-03-28 15:24:49

感谢您在我的一个问题中再次回复;） – Bernardo 2011-03-28 17:20:05

我能想到的一些解释的：

更改基线代码测量产生不利影响
无效的性能假设
身份不明的微-optimizations。
不现实的基准。

你说你的代码被简化了。我会尝试使用原始的基准代码，并看看会发生什么。如果数字更真实，则可以将原始基准代码与简化代码进行比较，以确定差异。

来源

2011-03-22 19:08:23

当我说简化它只是一个基本的数学问题（我刚刚加入了10^9和10^-3（转换为秒）的因素，没有别的）。至于你指出的解释，我会再次检查代码，看看我是否错过了一些东西 – Bernardo 2011-03-22 19:28:33

请注意，通过使用共享内存，纹理内存等，有时可能会超出理论带宽。这通常意味着您正在使用某些专用硬件支持的功能（如内置双线性纹理插值等），可能无意中。

除了罗伯特哈维提到的原因之外，厂商还可能出厂超频卡（尽管GeForce比Quadros更常见）。总体而言，如果您接近或超过理论带宽（无论是在内存还是计算中），我都会说你的表现不错。

来源

2011-03-23 13:02:43 peakxu

我是一个真正意义上的图形卡的小老虎，3个月前他们只是为了玩游戏。不过，我现在的结果是200 Gb/s（比理论值高4倍），并且我已经阅读了很多关于矩阵 - 矩阵乘法的讲座（从中我已经学习了代码），并且他们没有获得这些带宽值。所以我可能会错误计算这些值，但我真的找不到错误 – Bernardo 2011-03-23 15:07:47

CUDA理论带宽vs有效带宽

回答

相关问题