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有效帶寬

回答

相關問題