Czy jest możliwe równoległe obliczanie sumy w OpenCL?

Question

Jestem nowicjuszem w OpenCL. Rozumiem jednak podstawy C/C++ i OOP. Moje pytanie brzmi następująco: czy w jakiś sposób można równolegle uruchomić zadanie obliczania sumy? Czy to teoretycznie możliwe? Poniżej opiszę co próbowałem zrobić:

Zadanie jest, na przykład:

double* values = new double[1000]; //let's pretend it has some random values inside 
double sum = 0.0; 

for(int i = 0; i < 1000; i++) { 
    sum += values[i]; 
} 

Co próbowałem zrobić w kernelu OpenCL (i czuję, że jest źle, ponieważ może to uzyskuje dostęp do ta sama zmienna "suma" z różnych wątków/zadań w tym samym czasie):

__kernel void calculate2dim(__global float* vectors1dim, 
          __global float output, 
          const unsigned int count) { 
    int i = get_global_id(0); 
    output += vectors1dim[i]; 
} 

Ten kod jest niepoprawny. Będę bardzo wdzięczny, jeśli ktoś mi odpowie, jeśli teoretycznie możliwe jest równoległe wykonywanie takich zadań, a jeśli tak - to w jaki sposób!

Answer 1

Jeśli chcesz sumować wartości swojej macierzy w sposób równoległy, powinieneś upewnić się, że zredukować rywalizację o numer i upewnić się, że nie ma zależności danych między wątkami.

Zależności danych spowodują, że wątki będą musiały czekać na siebie nawzajem, tworząc rywalizację, której nie można uzyskać, aby uzyskać prawdziwą równoległość.

Jednym ze sposobów, w jaki można to zrobić, jest podzielenie macierzy na tablice N, z których każda zawiera jakiś podrozdział oryginalnej tablicy, a następnie wywołanie funkcji jądra OpenCL z każdą inną tablicą.

Na koniec, kiedy wszystkie jądra wykonały ciężką pracę, można po prostu podsumować wyniki każdej z tablic w jedną. Ta operacja może być łatwo wykonana przez procesor.

Kluczem jest brak zależności między obliczeniami wykonanymi w każdym jądrze, więc musisz odpowiednio podzielić dane i przetwarzanie.

Nie wiem, czy twoje dane mają jakiekolwiek rzeczywiste zależności od twojego pytania, ale to dla ciebie, aby dowiedzieć się.

Answer 2

Fragment kodu, który podałem w celach informacyjnych, powinien wykonać zadanie.

E.g. masz elementy i rozmiar grupy roboczej: WS = 64. Zakładam, że N jest wielokrotnością 2 * WS (jest to ważne, jedna grupa robocza oblicza sumę 2 * elementów WS). Następnie trzeba uruchomić określające jądra:

globalSizeX = 2*WS*(N/(2*WS)); 

Wskutek suma tablica będzie mieć sum cząstkowych 2 * WS elementów. (na przykład suma [1] - będzie zawierać sumę elementów, których indeksy są od 2 * WS do 4 * WS-1).

Jeśli twój plik globalSizeX jest 2 * WS lub mniejszy (co oznacza, że ​​masz tylko jedną grupę roboczą), oznacza to, że skończyłeś. Po prostu użyj suma [0] jako wynik. Jeśli nie - musisz powtórzyć procedurę, tym razem używając sumy jako tablicy wejściowej i wyjściowej do innej tablicy (utwórz 2 tablice i ping-pong między nimi). I tak dalej, dopóki nie będziesz miał tylko jednej grupy roboczej.

Wyszukaj również algorytmy równoległe Hilli Steele/Blelloch. This artykuł może być przydatny także

Oto rzeczywisty przykład:

__kernel void par_sum(__global unsigned int* input, __global unsigned int* sum) 
{ 
    int li = get_local_id(0); 
    int groupId = get_group_id(0); 

    __local int our_h[2 * get_group_size(0)]; 
    our_h[2*li + 0] = hist[2*get_group_size(0)*blockId + 2*li + 0]; 
    our_h[2*li + 1] = hist[2*get_group_size(0)*blockId + 2*li + 1]; 

    // sweep up 
    int width = 2; 
    int num_el = 2*get_group_size(0)/width; 
    int wby2 = width>>1; 

    for(int i = 2*BLK_SIZ>>1; i>0; i>>=1) 
    { 

     barrier(CLK_LOCL_MEM_FENCE); 

     if(li < num_el) 
     { 
      int idx = width*(li+1) - 1; 
      our_h[idx] = our_h[idx] + our_h[(idx - wby2)]; 
     } 

     width<<=1; 
     wby2 = width>>1; 
     num_el>>=1; 
    } 

     barrier(CLK_LOCL_MEM_FENCE); 

    // down-sweep 
    if(0 == li) 
     sum[groupId] = our_h[2*get_group_size(0)-1]; // save sum 
}