Michał Borek

Aktualnie zajmuję się stworzeniem wolumetrycznego efektu dymu. Zarówno do tego przedsięwzięcia, jak i do testów przecięć z typowymi prymitywami może być potrzebny test przecięcia promienia (raya) i cylindra.

Jak to zrobić?

Otóż, jeżeli mamy nasz promień w postaci

R = Origin + V * Direction

oraz równanie cylindra w postaci..

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 :)

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.