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.
Ż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:
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.