Michał Borek

Tech blog

OpenCL – język obliczeń równoległych – już jest!

5 grudnia br. została upubliczniona wersja 1.0 specyfikacji języka OpenCL (Open Computing Language), który ma być pomostem pomiędzy API poszczególnych technologii równoległego przetwarzania danych na GPU i CPU (między innymi w technologii NVidia CUDA, ATI Stream na procesorach Cell (tych z PS3 :)).

Khronos Group (twórca specyfikacji) chwali się tym, iż język ten ma zapewnić pełną niezależność języka od wykorzystanej technologii, a lista firm zaangażowanych w projekt, czyli: 3DLABS, Activision Blizzard, AMD, Apple, ARM, Barco, Broadcom, Codeplay, Electronic Arts, Ericsson, Freescale, HI, IBM, Intel, Imagination Technologies, Kestrel Institute, Motorola, Movidia, Nokia, NVIDIA, QNX, RapidMind, Samsung, Seaweed, Takumi, Texas Instruments i Umeå University uspokaja mnie, jeżeli chodzi o powodzenie tego projektu :)

Wcześniej zajmowałem się stricte technologią CUDA, ale z racji tego iż gdzieś na półce kurzy się moje PS3, to technologię tą zacznę wykorzystywać z jeszcze większą chęcią.

W najbliższym czasie postaram się opisać szerzej język OpenCL (jak tylko przeczytam specyfikację :)), a także sprawdzę na ile język ten jest wygodny, elastyczny i.. czy działa.

A jeżeli zadziała, może być ciekawie :)

CUDA – przykład (dodawanie wektorów)

Na początek taki szkolny przykład dodawanie do siebie dwóch tablic wektorów.

Dodawanie kolejnych par wektorów odbywa się na karcie graficznej. Odpowiada za to funkcja vecAdd z kwalifikatorem global. Kwalifikator ten oznacza, że funkcja może być wykonywana zarówno jako host (CPU) jak i device (GPU).

Prosty kod odpowiadający za dodawanie 2 wektorów

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
#include <stdio.h>
#include <cutil.h>
__global__ void vecAdd (float3 a[10], float3 b[10]) {
    int i = threadIdx.x;
    a[i].x += b[i].x;
    a[i].y += b[i].y;
    a[i].z += b[i].z;
}

int main(int argc, char** argv) {
    CUT_DEVICE_INIT(argc, argv); // inicjalizacja GPU
    float3 *a_h, *b_h, *a_d, *b_d, *result;
    int arraySize = 10;
    int sizeInBytes = arraySize * sizeof(float3);
    a_h = (float3*)malloc(sizeInBytes); // inicjalizacja tablic (na hoscie)
    b_h = (float3*)malloc(sizeInBytes);
    result = (float3*)malloc(sizeInBytes);
    CUDA_SAFE_CALL(cudaMalloc((void**)&a_d, sizeInBytes)); // inicjalizacja tablic na GPU
    CUDA_SAFE_CALL(cudaMalloc((void**)&b_d, sizeInBytes));
    for(int i = 0; i < arraySize; i++) { // przykładowe dane
    a_h[i].x = 100.0f + i;
    a_h[i].y = 200.0f + i;
    a_h[i].z = 300.0f + i;
    b_h[i].x = 50.0f + i;
    b_h[i].y = 50.0f + i;
    b_h[i].z = 50.0f + i;

    CUDA_SAFE_CALL(cudaMemcpy(a_d, a_h, sizeInBytes, cudaMemcpyHostToDevice));
    // kopiujemy tablicę z hosta do GPU
    CUDA_SAFE_CALL(cudaMemcpy(b_d, b_h, sizeInBytes, cudaMemcpyHostToDevice));
    vecAdd<<<1, arraySize>>>(a_d, b_d); // wykonujemy dodawanie
    CUDA_SAFE_CALL(cudaMemcpy(result, a_d, sizeInBytes, cudaMemcpyDeviceToHost));
    // kopiujemy z powrotem
    for(int i = 0; i < arraySize; i++){
        printf("Wartość %d.elementu %f, %f, %f \n", i, result[i].x,
        result[i].y, result[i].z); // wyswietlamy :)
    }
    CUT_EXIT(0, NULL); // koniec
}

Wydaje się proste :)

W przyszłości postaram się opisać po kolei wszystkie szczegóły.

Początek zabawy z CUDA

Żeby zacząć zabawę, należy kupić zabawkę.

Postanowiłem zaopatrzyć się w sprzęt, który pozwoli w pełni wykorzystać technologię CUDA do obliczeń związanych z ray tracingiem. Zdecydowałem się na Asusa GTX260, który nie miał zbyt wygórowanej ceny, ale oferował to co misie lubią najbardziej, czyli dużo “thread proccessors”. Sprzęcik wygląda mniej więcej tak:

Asus GTX 260

Jeżeli chodzi o mechanizm CUDA, początkowo (czego nie dotyczytałem wcześniej) okazało się, że funkcje global, device (czyli te uruchamiane na karcie graficznej), nie obsługują rekursji! Więc algorytm rekursywnych testów przecięć się nie sprawdzi. Może jednak uda się zaimplementować równoległe przeprowadzanie testów cieni, co w przypadku scen z dużą liczbą świateł, może dość mocno przyspieszyć generowanie obrazu.

W kolejnych postach postaram się zaprezentować przykładowe zastosowanie CUDA do szybkich obliczeń, a także powiedzieć coś więcej o samej technologii.