Grid wieloblokowy, jeden wątek oblicza wiele sąsiednich elementów macierzy
wynikowej, obliczenia przy wykorzystaniu pamięci współdzielonej bloku wątków.
__global__ void matrixMultiplyWithDownsampling ( float * C , float * B , float * A , int
MATRIX_SIZE , int skip ) {
...
__shared__ float Sds [ BLOCK_SIZE ][ BLOCK_SIZE ];
float sum [ BLOCK_SIZE ];
for ( int x = 0 ; x < skip ; x ++)
sum [ x ] = 0.0f;
for ( int z = 0 ; z < MATRIX_SIZE ; z += BLOCK_SIZE ){
for ( int x = 0 ; x < skip ; x ++)
Sds [ ty ][ tx + x ] = A [( row )* MATRIX_SIZE + ( tx + x ) + z ];
__syncthreads ();
for ( int x = 0 ; x < skip ; x ++)
for ( int i = 0 ; i < BLOCK_SIZE ; i ++)
sum [[ x ] += Sds [ ty ][ i ] * B [ i * MATRIX_SIZE + col +
( z )* MATRIX_SIZE + x ];
__syncthreads ();
}
for ( int x = 0 ; x < skip ; x ++)
C [( row )* MATRIX_SIZE + col + x ] = sum [ x ];
}
Trzeci wariant podobnie jak drugi wykorzystuje wieloblokowy grid o wymiarach
MATRIX_SIZE/BLOCK_SIZE, oraz korzysta z pamięci współdzielonej. W tej funkcji jednak możemy
przydzielić wątkom zadanie wyliczenia nie tylko jednego elementu macierzy, ale także sąsiednich.
Zmienna skip określa ile elementów macierzy wynikowej ma obliczyć każdy wątek. W
zaproponowanym przez nas rozwiązaniu każdy wątek oblicza elementy znajdujące się w tym samym
wierszu i sąsiadujących kolumnach. Obrazuje to poniższy rysunek (liczby w tabeli oznaczają numer
wątku)
1 1 2 2
3 3 4 4
5 5 6 6
7 7 8 8
Teraz CGMA ma wartość nie tylko zależną od wielkości bloku (tak samo jak w przypadku drugiej
wersji programu), ale również od wartości skip - w naszym przypadku wynoszącej dwa.
Podsumowując wartość CGMA dla powyższego kodu wyznacza się na podstawie danego wzoru:
skip*wielkość_bloku*liczba_operacji/liczba_odwołań_do_pamięci. Tak samo jak w pozostałych
wariantach tak i tutaj wartość liczby operacji i odwołań do pamięci wynosi dwa, tak więc
podsumowując CGMA = 2* wielkość bloku (gdzie 2 to wybrana przez nas wartość skip)
Teoretyczne zajętości procesora dla poszczególnych wielkości bloku wątków wynoszą:
Dla bloku 8x8: 16,7%, dla bloku 16x16: 58,3%, dla bloku 32x32: 66,7%