WSPÓLNE PODCIĄGI
Wniosek 7.2. Przyspieszenie w stosunku do wersji sekwencyjnej dla procesora CPU jest Θ
bwbh η2
. (7.14)
Rozważając teraz wyznaczanie kolejnych pudełek, można zauważyć, że w pierwszym uru-chomieniu kodu jądra, przetwarzane jest tylko 1 pudełko, w drugim uruuru-chomieniu kodu jądra – 2 pudełka itd. Maksymalna liczba pudełek, które mogą być przetwarzane w jednym uruchomie-niu kodu jądra to min(n0, m0). Złożoność czasowa pojedynczego wykonania kodu jądra zależy liniowo od maksymalnej liczby pudełek przydzielonych do jednego multiprocesora, tzn. e pu-dełek wyznaczanych jest w de/η2e krokach. Z tego, sumaryczna liczba kroków (dla wszystkich wykonań kodu jądra) jest
Θ(max(n0, m0)) +Θ
n0m0 min(n0, m0,η1)
=Θ
n0m0 min(n0, m0,η1)
. (7.15)
Jako że pojedyncze pudełko obliczane jest w czasie (7.14), sumaryczna złożoność czasowa obliczania pudełek jest
Θ
n0m0 min(n0, m0,η1)
×Θ bwbh
η2
=Θ
nm
min(n0, m0,η1)η2
. (7.16)
Łącznie istniejeΘ(n0m0) pudełek, przez co sumaryczny czas dostępu do pamięci globalnej jest Θ
n0m0bw η2
. (7.17)
Wniosek 7.1. Złożoność czasowa zaprezentowanego algorytmu wyznaczania podciągu LCS (LCS-DP-CUDA) jest sumą (7.16) i (7.17):
Θ
nm
min(n0, m0,η1)η2
+ nm η2bh
=Θ nm
η2min(bn
w,bm
h,η1, bh)
!
. (7.18)
Wniosek 7.2. Przyspieszenie w stosunku do wersji sekwencyjnej dla procesora CPU jest Θ
η2min
n bw,m
bh,η1, bh
. (7.19)
7.6.4. Algorytm równoległości bitowej
W algorytmie równoległości bitowej (BP) dla problemu LCS (podrozdz. 7.3.4) na etapie przetwarzania wstępnego wyznacza się wektory masek Yx. W wersji równoległej algorytmu wektory te wyznaczane są w procesorze CPU, a następnie przesyłane są do pamięci globalnej procesora GPU. Wektory te mogą być używane przez algorytm równoległy na kilka różnych sposobów. Najszybszym sposobem jest wykonanie ich kopii w pamięci wspólnej. Niestety, jest to możliwe tylko dla małych alfabetów, np. dla bh= 32 i 16 KB (4096 słów) pamięci wspólnej oczywistym limitem na rozmiar alfabetu jestσ= 4096/bh= 128. Pamięć wspólna używana jest jednak także do innych celów, m.in. do przechowywania fragmentów ciągów A i B, wymiany informacji o przeniesieniach przy dodawaniu wektorów bitowych. Ponadto, aby wykorzystać w pełni możliwości procesorów GPU, liczba bloków powinna być kilkakrotnie większa niż liczba multiprocesorów. Wszytko to powoduje, że maksymalny rozmiar alfabetu, dla którego możliwe jest wczytanie wszystkich wektorów masek do pamięci wspólnej, jest znacznie mniej-szy i w praktyce wynosi 20–32.
Alternatywą jest wykorzystywanie wektorów masek bitowych znajdujących się w pamięci globalnej. Niestety, przy takim podejściu opóźnienia spowodowane dostępem do tej pamięci będą bardzo duże i w dalszych badaniach nie brano tego wariantu pod uwagę.
Trzecie podejście jest połączeniem dwóch powyższych, w którym starano się połączyć zale-ty każdego z opisanych podejść przy jednoczesnej minimalizacji ich wad. Najszybszym sposo-bem dostępu do pamięci globalnej jest odczyt połączony, tzn. kolejne wątki operują na ko-lejnych słowach z pamięci globalnej. Niestety, wątki wyznaczające macierz muszą to robić metodą drugiej przekątnej, co powoduje, że każdy kolejny wątek może potrzebować fragmen-tu wektora maski, dotyczącego innego symbolu alfabefragmen-tu (takiego, jaki opisuje odpowiednią kolumnę macierzy). Dlatego też nie jest możliwe, aby dostęp do wektorów masek przez ko-lejne wątki był połączony. Wstępne eksperymenty pokazały, że dobrym kompromisem jest wykonywanie 128-bitowego dostępu do pamięci globalnej przez co czwarty wątek. Oznacza to, że fragment wektora masek dla 4 sąsiednich wątków (wg= 32) odczytywany jest w spo-sób połączony. Wartości te umieszczane są w tablicy pomocniczej znajdującej się w pamięci wspólnej, a każdy wątek pobiera z tej tablicy fragment wektora masek, kiedy go potrzebuje.
Podejście to redukuje liczbę dostępów do pamięci globalnej o czynnik 4 w stosunku do wa-riantu, w którym każdy wątek bezpośrednio czyta z pamięci globalnej, a przy tym nie ob-ciąża tak mocno pamięci wspólnej jak wariant z wykonaniem kopii całej tablicy wektorów masek.
Struktury danych znajdujące się w pamięci globalnej procesora GPU to:
• tablica n słów dla ciągu A,
• tablicaσdm/wge słów dla wektorów masek bitowych,
7.6. Algorytmy dla procesorów graficznych 119
• tablica n słów dla przeniesień z górnej krawędzi,
• tablica dm/wge słów dla wektora bitowego z lewej krawędzi.
Sumaryczna zajętość pamięci globalnej jest wobec tegoΘ(n +σdm/wge) słów.
Po wykonaniu obliczeń dla wszystkich pudełek do pamięci RAM transferowana jest zawar-tość prawej kolumny wszystkich pudełek z prawej krawędzi, dzięki czemu w procesorze CPU można wyznaczyć liczbę bitów o wartości0, która jest wynikiem działania algorytmu.
Część kodu jądra dla tego algorytmu pokazana jest na rys. 7.16. Parametrami wywołania funkcji są:
• block_x, block_y – znaczenie podobne jak w algorytmie DP,
• g_A – wskaźnik do tablicy zawierającej ciąg A,
• g_masks – wskaźnik do tablicy wektorów masek bitowych,
• g_lr_b – wskaźnik do tablicy zawierającej części wektora V (prawa krawędź sąsiedniego z lewej strony pudełka),
• g_tb_b – wskaźnik do tablicy zawierającej przeniesienia powstałe w trakcie obliczania górnego sąsiedniego pudełka,
• max_block_x – dodatkowa wartość pomocna do stwierdzenia, czy bieżące pudełko jest ostatnie w poziomie; jeśli tak, to zamiast prawej krawędzi zapisywana jest wg_lr_b liczba zer w V , co przyspiesza wyznaczenie końcowego wyniku.
Przedstawiony kod zawiera dwa warianty (o wyborze decyduje obecność makrodefinicji SMALL_ALPH). Jeśli alfabet jest mały, to wszystkie wektory masek bitowych wczytywane są do pamięci wspólnej (wiersze18–19), skąd są pobierane, kiedy jest to potrzebne (wiersz 32).
W przeciwnym przypadku wektory masek bitowych są pobierane z pamięci globalnej do tablicy pośredniejp_masks (wiersze 35–37) i są stamtąd odczytywane (wiersze 37, 39).
Analiza teoretycznego przyspieszenia algorytmu proponowanego w stosunku do algorytmu sekwencyjnego dla procesora CPU jest podobna do analizy przeprowadzonej dla algorytmu programowania dynamicznego, z tą różnicą, że wgwierszy przetwarzanych jest na raz, co daje całkowity czas wykonywanych obliczeń
Θ
nm
min(n0, m0,η1)η2wg
. (7.20)
Dostęp do pamięci wspólnej w obu wariantach algorytmu równoległości bitowej odbywa się bez konfliktów banków, a więc nie pojawia się tu serializacja wątków. Czas dostępu do pa-mięci globalnej jest jednak znacznie większy niż dla algorytmu opartego na programowaniu dynamicznym. W obu wersjach algorytmu równoległości bitowej (ze wstępnym wczytaniem wektorów masek oraz z wczytywaniem ich do tablicy pośredniej) występuje następująca liczba połączonych dostępów do pamięci globalnej dla każdego pudełka:Θ(bw/η2) (ładowanie części ciągu A oraz wczytywanie/zapisywanie górnej/dolnej krawędzi,Θ(bh/(η2wg)) (wczytywanie-/zapisywanie lewej i prawej krawędzi).
1 _ _ g l o b a l _ _ void Kernel_LCS_BP ( unsigned i n t block_x , unsigned i n t block_y , 2 unsigned i n t ∗g_A , unsigned i n t ∗g_masks ,
3 unsigned i n t ∗ g_lr_b , unsigned i n t ∗ g_tb_b , i n t max_block_x ) 4 {
5 __shared__ unsigned i n t sB [2∗BLK_HEIGHT ] ; 6 __shared__ unsigned i n t s_A [2∗BLK_HEIGHT ] ; 7 # i f d e f SMALL_ALPH
8 __shared__ unsigned i n t p_masks [ ALPH_SIZE ] [ BLK_HEIGHT ] ; 9 # e l s e
10 __shared__ unsigned i n t p_masks [ 4 ] [ BLK_HEIGHT ] ; 11 # e n d i f
12 unsigned i n t V, V2 , tB ;
13 c o n s t unsigned i n t n u m _ t h r e a d s = blockDim . x ; 14 c o n s t unsigned i n t t i d = t h r e a d I d x . x ;
15 u i n t 4 ∗ u4g_masks = ( u i n t 4 ∗) g_masks ; 16 . . .
17 # i f d e f SMALL_ALPH
18 f o r ( i = 0 ; i < ALPH_SIZE ; ++ i )
19 p_masks [ i ] [ t i d ] = g_masks [ ( i << SHIFT_MASKS ) + o r i g _ y + t i d ] ; 20 # e n d i f
21 . . .
22 f o r ( ; i < BLK_WIDTH; ) { 23 i n t i_max = i + n u m _ t h r e a d s ;
24 load_A ( s_A , g_A , o r i g _ x , t i d , i , o r i g _ y ) ; 25 l o a d _ s B ( sB , g_t b_b , o r i g _ x , t i d , i ) ; 26 _ _ s y n c t h r e a d s ( ) ;
27
28 f o r ( ; i < i_max ; ++ i ) {
29 i n t n _ i = ( i − t i d ) & BLK_MASK;
30 unsigned i n t m;
31 # i f d e f SMALL_ALPH
32 m = p_masks [ s_A [ n _ i ] ] [ t i d ] ; 33 # e l s e
34 i f ( ( t i d & 3 ) == 0 ) {
35 u i n t 4 mm = u4g_masks [ s_A [ n _ i ] + ( t i d >> 2 ) ] ;
36 p_masks [ n _ i & 3 ] [ t i d +1] = mm. y ; p_masks [ n _ i & 3 ] [ t i d +2] = mm. z ; 37 p_masks [ n _ i & 3 ] [ t i d +3] = mm. w; m = mm. x ;
38 }
39 e l s e m = p_masks [ n _ i & 3 ] [ t i d ] ; 40 # e n d i f
41 tB = V & m; V2 = V + sB [ n _ i ] ; V2 += tB ;
42 sB [ n _ i ] = V2 < V; V = ( V2 | (V − tB ) ) ; 43 _ _ s y n c t h r e a d s ( ) ;
44 }
45 g _ t b _ b [ o r i g _ x + t i d + i −2∗BLK_HEIGHT ] = sB [ ( t i d +i −2∗BLK_HEIGHT ) & BLK_MASK ] ;
46 }
47 . . .
48 i f ( max_block_x != block_x −b l o c k I d x . x ) g _ l r _ b [ o r i g _ y + t i d ] = V;
49 e l s e g _ l r _ b [ o r i g _ y + t i d ] = __popc (~V ) ;
50 }
Rys. 7.16. Część programu jądra dla algorytmu LCS-BP-CUDA rozwiązującego problem LCS metodą równoległości bitowej w procesorze GPU
Fig. 7.16. A part of the kernel code for the LCS BP algorithm (LCS-BP-CUDA) at GPU
7.6. Algorytmy dla procesorów graficznych 121 Wersja ze wstępnym wczytaniem wektorów masek
W wersji ze wstępnym wczytaniem wektorów masek występuje konieczność wczytania ca-łych wektorów masek dla części ciągu B, co wymagaΘ(σbh/(η2wg)) połączonych dostępów do pamięci. Ponieważσ≤ bw(jeśli to nie jest spełnione, to wystarczy wczytać te wektory ma-sek, dla których występują symbole w odpowiedniej części ciągu A), całkowity czas dostępów do pamięci globalnej jest
Θ Złożoność czasowa tego algorytmu jest sumą (7.20) oraz (7.21):
Θ nm
Złożoność czasowa wersji sekwencyjnej dla procesora CPU jestΘ(nm/wc), a więc:
Wniosek 7.3. Przyspieszenie wersji ze wstępnym wczytaniem wektorów masek w stosunku do