come ottimizzare la moltiplicazione della matrice usando OpenACC?

Question

Sto imparando OpenACC (con il compilatore di PGI) e sto cercando di ottimizzare l'esempio di moltiplicazione della matrice. L'implementazione più veloce mi è venuta finora è la seguente:

void matrix_mul(float *restrict r, float *a, float *b, int N, int accelerate){ 

#pragma acc data copyin (a[0: N * N ], b[0: N * N]) copyout (r [0: N * N ]) if(accelerate) 
{ 
# pragma acc region if(accelerate) 
{ 
# pragma acc loop independent vector(32) 
for (int j = 0; j < N; j ++) 
{  
    # pragma acc loop independent vector(32) 
    for (int i = 0; i < N ; i ++) 
    { 
     float sum = 0; 
     for (int k = 0; k < N ; k ++) { 
     sum += a [ i + k*N ] * b [ k + j * N ]; 
     } 
     r[i + j * N ] = sum ; 
    } 
} 
} 
} 

Questo si traduce in blocchi di filettatura di dimensioni 32x32 discussioni e mi dà le migliori prestazioni finora. Qui sono i parametri di riferimento:

Matrix multiplication (1500x1500): 
GPU: Geforce GT650 M, 64-bit Linux 

Data sz    : 1500  
Unaccelerated: 
    matrix_mul() time : 5873.255333 msec 
Accelerated: 
    matrix_mul() time : 420.414700 msec 

Data size    : 1750 x 1750  
    matrix_mul() time : 876.271200 msec 
Data size    : 2000 x 2000  
    matrix_mul() time : 1147.783400 msec 
Data size    : 2250 x 2250  
    matrix_mul() time : 1863.458100 msec 
Data size    : 2500 x 2500  
    matrix_mul() time : 2516.493200 msec 

Purtroppo ho capito che il codice CUDA generato è molto primitivo (ad esempio neppure non utilizza la memoria condivisa) e quindi non può competere con il programma CUDA ottimizzato a mano. Come un'implementazione di riferimento ho preso Arrayfire lib con i seguenti risultati:

Arrayfire 1500 x 1500 matrix mul 
CUDA toolkit 4.2, driver 295.59 
GPU0 GeForce GT 650M, 2048 MB, Compute 3.0 (single,double) 
Memory Usage: 1932 MB free (2048 MB total) 
af: 0.03166 seconds 

Arrayfire 1750 x 1750 matrix mul 
af: 0.05042 seconds 
Arrayfire 2000 x 2000 matrix mul 
af: 0.07493 seconds 
Arrayfire 2250 x 2250 matrix mul 
af: 0.10786 seconds 
Arrayfire 2500 x 2500 matrix mul 
af: 0.14795 seconds 

mi chiedo se ci sono suggerimenti su come ottenere prestazioni migliori da OpenACC? Forse la mia scelta delle direttive non è giusta?

Answer 1

Stai arrivando a un errore di 14 volte, che è abbastanza buono per il compilatore di PGI in base alla mia esperienza.

Prima di tutto, stai compilando con -Minfo? Ciò ti darà un sacco di feedback dal compilatore riguardo alle scelte di ottimizzazione.

Si sta utilizzando un blocco di thread 32x32, ma nella mia esperienza i blocchi di thread 16x16 tendono a ottenere prestazioni migliori. Se ometti le clausole vettoriali (32), quale scheduling sceglie il compilatore?

Dichiarare a e b con restrizioni potrebbe consentire al compilatore di generare codice migliore.

Solo guardando il codice, non sono sicuro che la memoria condivisa possa aiutare le prestazioni. La memoria condivisa aiuta solo a migliorare le prestazioni se il codice può memorizzare e riutilizzare i valori lì invece di andare alla memoria globale. In questo caso non stai riutilizzando nessuna parte di aob dopo averlo letto.

Vale anche la pena notare che ho avuto brutte esperienze con il compilatore di PGI quando si tratta di utilizzo della memoria condivisa. A volte fa cose buffe e memorizza i valori sbagliati (sembra accadere soprattutto se si itera un ciclo all'indietro), generando risultati errati. In realtà devo compilare la mia attuale applicazione usando l'opzione non documentata -ta = nvidia, nocache per farlo funzionare correttamente, ignorando del tutto l'uso della memoria condivisa.