jak zoptymalizować mnożenie macierzy za pomocą OpenACC?

Question

Uczę się OpenACC (z kompilatorem PGI) i próbuję zoptymalizować przykład mnożenia macierzy. Najszybsza realizacja wymyśliłem tak daleko jest następujące:

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 ; 
    } 
} 
} 
} 

Wynika to w blokach o rozmiarze 32x32 nici nici i daje mi najlepszą wydajność do tej pory. Oto benchmarki:

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 

Niestety zdałem sobie sprawę, że wygenerowany kod CUDA jest dość prymitywne (np nawet nie korzystają z pamięci współdzielonej), a zatem nie może konkurować z ręcznie zoptymalizowany programu CUDA. W implementacji referencyjnej Wziąłem Arrayfire lib z następującymi wynikami:

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 

Zastanawiam się, czy istnieją jakieś sugestie, jak uzyskać lepszą wydajność od OpenACC? Być może mój wybór dyrektyw nie jest właściwy?

Answer 1

Masz prawo przyspieszyć 14x, co jest całkiem dobre dla mojego kompilatora PGI.

Po pierwsze, czy kompilujesz za pomocą -Minfo? To da ci wiele informacji zwrotnych od kompilatora dotyczących wyboru optymalizacji.

Używasz bloku gwintów 32x32, ale z mojego doświadczenia wynika, że ​​bloki gwintów 16x16 mają tendencję do lepszej wydajności. Jeśli pominiesz klauzule vector (32), jakie harmonogramy wybiera kompilator?

Zgłaszanie a i b z ograniczeniem może pozwolić kompilatorowi wygenerować lepszy kod.

Tylko patrząc na twój kod, nie jestem pewien, czy pamięć współdzielona pomogłaby wydajności. Pamięć współdzielona pomaga tylko poprawić wydajność, jeśli twój kod może przechowywać i ponownie używać wartości zamiast przechodzić do pamięci globalnej. W takim przypadku nie używasz żadnej części a lub b po jej przeczytaniu.

Warto również zauważyć, że miałem złe doświadczenia z kompilatorem PGI, jeśli chodzi o wykorzystanie pamięci współdzielonej. Czasami robi zabawne rzeczy i buforuje błędne wartości (wydaje się, że dzieje się tak głównie, jeśli przerwiesz pętlę wstecz), generując złe wyniki. Muszę teraz skompilować moją bieżącą aplikację przy użyciu nieudokumentowanej opcji -ta = nvidia, nocache, aby uzyskać poprawne działanie, pomijając użycie pamięci współdzielonej.