Michał Borek

Tech blog

Rozproszony Ray Tracing

Na zajęcia z systemów rozproszonych postanowiłem rozwinąć lekko swój projekt ray tracera, tworzony w zeszłym roku.

Projekt oparty został o bibliotekę QT w wersji 4.5. Komunikacja odbywa się za pomocą socketów TCP.

Architektura tego rozwiązania opiera się o 3 elementy:

Klient – program, w którym operuje użytkownik chcący wygenerować obraz metoda ray tracing.
Serwer zarządzający – który rozsyła części obrazu do generowania przez serwery robocze
Serwery robocze – pracujące nad generowaniem poszczególnych części obrazu.

Jak to bywa w projektach studenckich, także i ten jest napisany obecnie dość chaotycznie i zawiera na pewno wiele błędów. Będą one jednak usuwane z biegiem czasu, a w chwili obecnej istotne jest raczej to, że projekt mniej więcej działa 🙂

Na początek jeden screen klienta:

Rozproszony Ray Tracing

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.