Algorytm programowania dynamicznego

Przekształcenie sekwencyjnego algorytmu programowania dynamicznego (DP) dla proble-mu LCS w algorytm równoległy dla procesorów GPU, zgodnie ze schematem zaproponowa-nym powyżej, jest stosunkowo proste. Przetwarzanie wstępne (rys. 7.13, wiersz 2) nie jest tu wykonywane. Struktury danych znajdujące się w pamięci globalnej procesora GPU to:

• odpowiednio n oraz m słów dla ciągów A oraz B,

• tablica n słów dla komórek z górnej krawędzi (b^wdla każdego równolegle przetwarzanego pudełka),

• tablica m słów dla komórek z lewej krawędzi (b^h dla każdego równolegle przetwarzanego pudełka),

• tablica 3m/b^h słów dla komórek zawierających wartość prawej dolnej komórki lewego-górnego sąsiedniego pudełka do pudełka bieżącego.

7.6. Algorytmy dla procesorów graficznych 115 Z powyższego wynika, że całkowita zajętość pamięci globalnej w tym algorytmie jestΘ(n+

m) =Θ(n) słów. Pamięć dla lewych i górnych krawędzi może być liniowa dzięki temu, że każde równolegle obliczane pudełko wykorzystuje inny obszar tych tablic. Wymaganych jest 3 razy więcej słów w tablicy przechowującej wartości prawych dolnych komórek pudełek, ponieważ wartości z dwóch ostatnio przetwarzanych pudełek z bieżącego wiersza muszą być znane, a tak-że gdzieś musi być możliwość przechowania wartości analogicznej komórki z bieżącego pudeł-ka. Wynik końcowy znajduje się w prawej dolnej komórce prawego dolnego pudełka, a więc tylko ta liczba musi być skopiowana do pamięci RAM i żadne inne obliczenia nie są konieczne.

Część kodu tego algorytmu pokazana jest na rys. 7.15. Parametrami wywołania funkcji są:

• block_x – maksymalny indeks poziomy pudełka wyznaczanego w bieżącym wykonaniu jądra,

• block_y – indeks pionowy pudełka wskazanego powyżej,

• g_A – wskaźnik do tablicy (umieszczonej w pamięci globalnej) rozmiaru n zawierającej w i-tej komórce wartość ai oraz M(i, j⁰), gdzie j⁰ jest największym takim indeksem j, dla którego wartość M(i, j) została już obliczona,

• g_B – wskaźnik do tablicy (umieszczonej w pamięci globalnej) rozmiaru m zawierającej w j-tej komórce wartość bjoraz M(i⁰, j), gdzie i⁰ jest największym takim indeksem i, dla którego wartość M(i, j) została już obliczona,

• g_res – wskaźnik do tablicy (umieszczonej w pamięci globalnej) rozmiaru 3dm/we zawie-rającej wartości prawych dolnych komórek ostatnio obliczanych pudełek.

Symbole ciągu B wczytywane są z pamięci globalnej do rejestrów sB (wiersze 17–28).

Ciąg A wczytywany jest do tablicys_A (znajdującej się w pamięci wspólnej) w partiach (np.

wiersze 15–16, 28–29), ponieważ s_A funkcjonuje jako bufor cykliczny. Za każdym razem, kiedy wątek oblicza wartość pojedynczej komórki, znane są dla niej wartości:

• komórki będącej górnym sąsiadem – wartość w tablicy s_u (z pamięci wspólnej),

• komórki będącej lewym sąsiadem – wartość w zmiennej rejestrowej v,

• komórki będącej lewym-górnym sąsiadem – wartość w zmiennej rejestrowej v_upp.

Opierając się na tych wartościach oraz na symbolach ciągów A, B, wyznaczana jest wartość bieżącej komórki, która jest następnie umieszczana w tablicy s_u (wiersze 33–34). Za każ-dym razem kiedy wyznaczona jest połowa wartości znajdujących się w tablicys_u, wartości te kopiowane są do tablicyg_A umieszczone w pamięci globalnej (wiersze 38, 44). Po zakończe-niu przetwarzania pudełka wartości komórek z prawej krawędzi są zapisywane do tablicyg_B (linia 43). Ponadto, wartość prawej-dolnej komórki jest umieszczana w tablicyg_res.

Dla przejrzystości prezentacji, pokazany kod dotyczy tylko wyznaczania środkowej części pudełka, tj. odpowiadającej etapom 61–64oraz 91–94 z rys. 7.14. Kod dla etapów 41–44 oraz 11₁–113jest podobny, z tą różnicą, że niektóre wątki nie wykonują obliczeń.

1 _ _ g l o b a l _ _ void Kernel_LCS_DP ( c o n s t unsigned i n t block_x ,

2 c o n s t unsigned i n t block_y , u i n t 2 ∗g_A , u i n t 2 ∗g_B , unsigned i n t ∗ g _ r e s ) 3 {

4 __shared__ unsigned i n t s_A [2∗BLK_HEIGHT ] ; 5 __shared__ unsigned i n t s_u [2∗BLK_HEIGHT ] ; 6 unsigned i n t sB ;

7 i n t i ;

8 unsigned i n t v , v_upp ; 9

10 c o n s t unsigned i n t n u m _ t h r e a d s = BLK_HEIGHT ; 11 c o n s t unsigned i n t t i d = t h r e a d I d x . x ;

12 c o n s t unsigned i n t bx0 = b l o c k _ x − b l o c k I d x . x ; 13 c o n s t unsigned i n t by0 = b l o c k _ y + b l o c k I d x . x ; 14 . . .

15 u i n t 2 tmp = g_A [ o r i g _ x + i + t i d ] ;

16 s_A [ ( t i d + i ) & BLK_MASK] = tmp . x ; s_u [ ( t i d + i ) & BLK_MASK] = tmp . y ; 17 tmp = g_B [ o r i g _ y + t i d ] ;

18 sB = tmp . x ; v = tmp . y ; 19

20 / / L e f t−upper

21 i f ( t i d > 0 ) v_upp = g_B [ o r i g _ y + t i d − 1 ] . y ; 22 e l s e i f ( o r i g _ y == 0 ) v_upp = 0 ;

23 e l s e v_upp = g _ r e s [ o f f _ r e s _ r e a d + by0 − 1 ] ; 24 . . .

25

26 f o r ( ; i < BLK_WIDTH; ) { 27 i n t i_max = i + n u m _ t h r e a d s ; 28 u i n t 2 tmp = g_A [ o r i g _ x + i + t i d ] ;

29 s_A [ ( t i d + i ) & BLK_MASK] = tmp . x ; s_u [ ( t i d + i ) & BLK_MASK] = tmp . y ; 30

31 f o r ( ; i < i_max ; ++ i ) {

32 i n t n _ i = ( i − t i d ) & BLK_MASK;

33 v = max ( v , v_upp + ( sB == s_A [ n _ i ] ) ) ; v_upp = s_u [ n _ i ] ;

34 v = max ( v_upp , v ) ; s_u [ n _ i ] = v ;

35 _ _ s y n c t h r e a d s ( ) ;

36 }

37

38 g_A [ o r i g _ x + t i d + i −2∗ num_threads ] . y = s_u [ ( t i d +i −2∗ num_threads ) & BLK_MASK ] ; 39 _ _ s y n c t h r e a d s ( ) ;

40 }

41 . . . 42

43 g_B [ o r i g _ y + t i d ] . y = v ;

44 g_A [ o r i g _ x + t i d + i −2∗ num_threads ] . y = s_u [ ( t i d + i −2∗ num_threads ) & BLK_MASK ] ; 45

46 i f ( t i d == BLK_HEIGHT − 1) 47 g _ r e s [ o f f _ r e s _ w r i t e +by0 ] = v ; 48

49 _ _ s y n c t h r e a d s ( ) ; 50 }

Rys. 7.15. Część kodu jądra algorytmu LCS-DP-CUDA rozwiązującego problem LCS metodą DP w pro-cesorze GPU

Fig. 7.15. A part of the kernel code for the LCS DP algorithm (LCS-DP-CUDA) at GPU

Poniżej rozważone zostanie teoretyczne przyspieszenie tej wersji algorytmu w stosunku do algorytmu sekwencyjnego, działającego w czasieΘ(nm). Do obliczenia pojedynczego pudełka potrzebaΘ(bw/η2) połączonych (ang. coalesced) dostępów do pamięci globalnej, aby wczytać związaną z tym pudełkiem część ciągu A. Podobnie potrzebaΘ(bh/η2) połączonych dostępów do pamięci globalnej, aby wczytać części ciągu B. Taka sama liczba połączonych dostępów do pamięci globalnej jest wymagana, aby wczytać oraz zapisać wartości na krawędziach

pu-7.6. Algorytmy dla procesorów graficznych 117 dełka. Z tego wynika, że sumaryczny czas dostępu do pamięci globalnej dla jednego pudełka jestΘ((bw+ bh)/η2) =Θ(bw/η2). Dostępy do pamięci globalnej przez różne multiprocesory wykonywane są sekwencyjnie, co zostanie wzięte pod uwagę nieco później.

Każdy multiprocesor wyposażony jest w osobną pamięć wspólną i nie występuje serializacja wątków operujących na pamięci wspólnej, jeśli należą one do różnych bloków. Wszystkie do-stępy do pamięci wspólnej wykonywane są bez konfliktów banków (ang. bank conflicts), a więc równolegle. Liczba przebiegów pętli w ramach jednego pudełka jestΘ(bw+bh) =Θ(bw), z cze-go liczba operacji wykonywanych dla wyliczenia jednecze-go pudełka to:

Θ

b_wb_h η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(n⁰, m⁰). 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(n⁰, m⁰)) +Θ

n⁰m⁰ min(n⁰, m⁰,η1)



=Θ

n⁰m⁰ min(n⁰, m⁰,η1)



. (7.15)

Jako że pojedyncze pudełko obliczane jest w czasie (7.14), sumaryczna złożoność czasowa obliczania pudełek jest

Θ

n⁰m⁰ min(n⁰, m⁰,η1)



×Θ b_wb_h

η2



=Θ

nm

min(n⁰, m⁰,η1)η2



. (7.16)

Łącznie istniejeΘ(n⁰m⁰) pudełek, przez co sumaryczny czas dostępu do pamięci globalnej jest Θ

n⁰m⁰b_w η2



. (7.17)

Wniosek 7.1. Złożoność czasowa zaprezentowanego algorytmu wyznaczania podciągu LCS