opencl
Podstawy jądra

Wprowadzenie

Jądro w skali szarości

Zbudujmy jądro, aby wygenerować obraz w skali szarości. Użyjemy danych obrazu, które są zdefiniowane za pomocą uint dla każdego komponentu i na zamówienie RGBA.

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

Teraz przejdźmy krok po kroku przez ten kod. Pierwszy wiersz tworzy zmienną w ciągłym obszarze pamięci typu sampler_t. Ten sampler służy do dalszego określania dostępu do naszych danych obrazu. Pełna dokumentacja znajduje się w dokumentach Khronos.

Przydzieliliśmy wejście jako read_only, a wyjście jako write_only, zanim wywołaliśmy nasze jądro, więc dodajemy tutaj te modyfikatory.

image2d i image3d są zawsze przydzielane w pamięci globalnej, dlatego tutaj możemy pominąć modyfikator __global.

Następnie otrzymujemy nasz identyfikator wątku, który określa piksel, który zamierzamy przekonwertować na skalę szarości. Sprawdzamy również rozmiar, aby upewnić się, że nasz wątek nie ma dostępu do nieprzydzielonej pamięci. Spowoduje to awarię jądra, jeśli o tym zapomnisz.

Po upewnieniu się, że jesteśmy legalnym wątkiem, odczytujemy piksel z naszego obrazu wejściowego. Następnie przekształcamy go na zmiennoprzecinkowy, aby uniknąć utraty miejsc dziesiętnych, wykonujemy obliczenia, konwertujemy go z powrotem i zapisujemy na wyjściu.

Jądro Skelleton

Przejdźmy przez najprostsze dostępne jądro i niektóre jego odmiany

__kernel void myKernel() {
}

Jądro, które można uruchomić z głównego kodu, jest identyfikowane przez słowo kluczowe __kernel. Funkcja jądra może mieć tylko typ zwracany void.

__kernel void myKernel(float a, uint b, byte c) {

}

Oczywiście możesz tworzyć więcej funkcji, które nie są ujawniane jako jądra. W takim przypadku możesz po prostu pominąć modyfikator __kernel.

Funkcja może ujawniać zmienne, jak każda inna funkcja C / C ++. Jedyną różnicą jest to, kiedy chcesz odwoływać się do pamięci. Dotyczy to wszystkich wskaźników, bez względu na to, czy są argumentami, czy są używane w kodzie.

float*  ptr;

jest wskaźnikiem do obszaru pamięci, do którego dostęp ma tylko wykonujący wątek. W rzeczywistości jest to to samo co

__private float* ptr;

Dostępne są cztery różne modyfikatory regionu pamięci. W jądrze zwykle nie musisz się tym martwić, ale w przypadku argumentów jest to niezbędne.

• __global: ten modyfikator odnosi się do wskaźnika umieszczonego w pamięci globalnej
• __constant: odnosi się do stałego wskaźnika pamięci
• __local: odnosi się do wskaźnika pamięci współdzielonej
• __private: odnosi się do lokalnego wskaźnika pamięci

Ponadto możemy zdefiniować, w jaki sposób chcemy uzyskać dostęp do pamięci

• bez modyfikatora: odczyt i zapis
• __tylko czytać
• __zapisz_tylko

Te flagi muszą pasować do sposobu, w jaki przydzieliliśmy bufor pamięci na naszym hoście.

ID jądra

Aby poprawnie pracować z danymi, każdy wątek musi znać swoją pozycję w puli wątków / globalnej puli wątków. Można to archiwizować za pomocą

get_local_id($dim);
get_global_id($dim);

Te dwie funkcje zwracają pozycję wątku względem bloku wątków lub wszystkich wątków.

get_working_dim();

Pobiera całkowitą liczbę wymiarów, z którymi jądro zostało uruchomione.

get_local_size($dim);
get_global_size($dim);

Pobiera całkowitą liczbę wątków w bloku wątków lub łącznie dla danego wymiaru.

Uwaga: upewnij się, że Twój wątek nie przekracza rozmiaru danych. Jest to bardzo prawdopodobne i zawsze należy to sprawdzić.

Wektory w OpenCL

Każdy podstawowy typ opencl ma wersję wektorową. Możesz użyć typu wektora, dodając liczbę pożądanych komponentów po typie. Obsługiwana liczba komponentów to 2,3,4,8 i 16. OpenCL 1.0 nie oferuje trzech komponentów.

Możesz zainicjalizować dowolny wektor na dwa sposoby:

• Podaj pojedynczy skalar
• Spełnij wszystkie elementy

float4 a = (float4)(1); //a = (1, 1, 1, 1)

lub

float4 b = (float4)(1, 2, 3, 4);
float4 c = (float4)(1, (float3)(2));

lub dowolna inna kombinacja wektorów, które spełniają liczbę składników. Aby uzyskać dostęp do elementów wektora, możesz użyć różnych metod. Możesz albo użyć indeksowania:

a[0] = 2;

lub użyj literałów. Zaletą literałów jest to, że możesz je łączyć w dowolny sposób, zrób to za chwilę. Możesz uzyskać dostęp do wszystkich komponentów wektorowych za pomocą

a.s0 = 2; // same as a[0] = 2

możesz także łączyć wiele komponentów w nowy wektor

a.s02 = (float2)(0, 0); // same as  a[0] = 0; a[2] = 0; or even a.s20 = (float2)(0, 0)

możesz zmienić kolejność komponentów w dowolny sposób.

a.s1423 = a.s4132; // flip the vector

ale nie możesz zrobić czegoś takiego

a.s11 = ... // twice the same component is not possible

Istnieje kilka wygodnych skrótów dostępu do komponentów wektorowych. Poniższe skróty dotyczą tylko rozmiarów 2, 4, 8 i 16

a.hi //=a.s23 for vectors of size 4, a.4567 for size 8 and so on.
a.lo //=a.s01
a.even //=a.s02
a.odd //=a.13

Dla rozmiarów wektorów 2,3 i 4 istnieje kilka dodatkowych skrótów

a.x //=a.s0
a.y //=a.s1
a.z //=a.s2
a.w //=a.s3

Jądro korekcji gamma

Spójrzmy na jądro korekcji gamma

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Gamma(__read_only image2d_t input, __write_only image2d_t output, __constant float gamma) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color = pow(color, (float4)(gamma));
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

Teraz przejdźmy krok po kroku przez ten kod. Pierwszy wiersz tworzy zmienną w ciągłym obszarze pamięci typu sampler_t. Ten sampler służy do dalszego określania dostępu do naszych danych obrazu. Pełna dokumentacja znajduje się w dokumentach Khronos.

Przydzieliliśmy wejście jako read_only, a wyjście jako write_only, zanim wywołaliśmy nasze jądro, więc dodajemy tutaj te modyfikatory.

image2d i image3d są zawsze przydzielane w pamięci globalnej, dlatego tutaj możemy pominąć modyfikator __global. Nasza wartość gamma znajduje się w __stałej pamięci, więc również to określamy.

Następnie otrzymujemy nasz identyfikator wątku, który określa piksel, który będziemy poprawiać gamma. Sprawdzamy również rozmiar, aby upewnić się, że nasz wątek nie ma dostępu do nieprzydzielonej pamięci. Spowoduje to awarię jądra, jeśli o tym zapomnisz.

Po upewnieniu się, że jesteśmy legalnym wątkiem, odczytujemy piksel z naszego obrazu wejściowego. Następnie przekształcamy go na zmiennoprzecinkowy, aby uniknąć utraty miejsc dziesiętnych, wykonujemy obliczenia, konwertujemy go z powrotem i zapisujemy na wyjściu.