CUDA bande passante théorique vs bande passante effective

Question

J'ai un noyau CUDA qui multiplie deux matrices dont la largeur et la hauteur sont des multiples de la taille de bloc que j'utilise.

La Nvidia Quadro Fx 3800 J'utilise a une bande passante théorique de 50 Gb/s et je vais avoir des résultats étranges (efficace bande passante plus large que la bande passante théorique)

Je signalerai ici quelques résultats:

Avec Blocksize 2

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

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

Avec Blocksize 64

[1000] -> BW = 45072,12 Gb/s

Je pris la formule de la bande passante effective du Guide des meilleures pratiques Nvidia (je l'ai simplifié, mais son équivalent (à moins d'une erreur stupide)). Je pense que le noyau est très bien car il est très similaire (sinon égal) à certaines conférences Nvidia que j'ai lues et aussi parce que cela fonctionne correctement (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); 
    } 

Merci à l'avance

Answer 1

Je pense que le noyau est en train d'échouer silencieusement.

1. Avez-vous vérifié des erreurs après l'invocation du noyau ?

2. Le code fonctionne-t-il? Quels résultats avez-vous sur les timings ?

Answer 2

Je peux penser à un certain nombre d'explications:

1. Les modifications apportées au code de base qui affectent négativement les mesures
2. hypothèses de performance non valides
3. non identifiée micro -optimisations.
4. Repères irréalistes.

Vous dites que votre code est simplifié. Je voudrais essayer d'utiliser le code de référence d'origine, et voir ce qui se passe. Si les chiffres sont plus réalistes, vous pouvez comparer le code de référence d'origine avec votre code simplifié pour identifier les différences.

Answer 3

Notez qu'en utilisant la mémoire partagée, la mémoire de texture, etc., il est parfois possible de dépasser la bande passante théorique. Cela signifie souvent que vous utilisez certaines fonctions matérielles dédiées (telles que l'interpolation de texture bilinéaire intégrée, etc.), peut-être involontairement. Outre les raisons évoquées par Robert Harvey, il existe également un risque d'overclocking de cartes par les fournisseurs (bien que plus fréquent avec GeForce qu'avec Quadros). Dans l'ensemble, je dirais que vous vous débrouillez bien si vous vous rapprochez de la bande passante théorique ou que vous la dépassez (en mémoire ou en calcul).