库达理论带宽与有效带宽

Question

我有一个CUDA内核，它乘以两个矩阵，它们的宽度和高度是我使用的块大小的倍数。

我正在使用的Nvidia Quadro Fx 3800的理论带宽为50 Gb/s，并且我有一些奇怪的结果(有效带宽大于理论带宽)。

我将在这里发布一些结果：

与细条2

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

带小块64的

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

我采用了Nvidia最佳实践指南中的有效带宽公式(我已经简化了它，但它的等效性(除非有一个愚蠢的错误))。我认为内核很好，因为它与我读过的一些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);
    }

提前感谢

Answer 1

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

1. 您在内核调用之后检查了错误吗？
2. 代码工作吗？
3. ，您对时间有什么结果?

Answer 2

请注意，通过使用共享内存、纹理内存等，有时可能会超出理论带宽。这通常意味着您正在使用一些专用的硬件支持的函数(例如内置的双线性纹理插值等)，可能是无意的。

除了罗伯特·哈维提到的原因外，厂商们还可能会在工厂里大打出手(尽管GeForce比Quadros更常见)。

总的来说，如果您接近或超过理论带宽(无论是在内存中还是在计算中)，我会说您做得很好。

Answer 3

我能想到一些解释：

measurements

• Invalid性能assumptions

• Unidentified micro-optimizations.

• Unrealistic基准测试基准代码的

1. 更改。

。

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