opencl
Kernel Basics

Invoering

Grijswaarden kernel

Laat een kernel bouwen om een grijswaardenafbeelding te genereren. We gebruiken afbeeldingsgegevens die worden gedefinieerd met behulp van uints voor elk onderdeel en bij bestelling 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);
    }
}

Laten we nu stap voor stap door die code lopen. De eerste regel maakt een variabele in het __constante geheugengebied van het type sampler_t. Deze sampler wordt gebruikt om de toegang tot onze afbeeldingsgegevens nader te specificeren. Raadpleeg de Khronos-documenten voor een volledige documentatie.

We hebben de input toegewezen als read_only en de output als write_only voordat we onze kernel aanroerden, dus we voegen die modificatoren hier toe.

image2d en image3d worden altijd toegewezen aan het wereldwijde geheugen, daarom kunnen we de __global-modifier hier weglaten.

Vervolgens krijgen we onze thread-ID die de pixel definieert die we naar grijswaarden gaan converteren. We vragen ook naar de grootte om te controleren of onze thread geen toegang krijgt tot niet-toegewezen geheugen. Dit zal je kernel zeker laten crashen als je dat vergeet.

Nadat we ervoor hebben gezorgd dat we een legitieme thread zijn, lezen we onze pixel uit onze invoerafbeelding. We converteren het vervolgens naar zwevend om verlies van decimalen te voorkomen, voeren enkele berekeningen uit, zetten het terug en schrijven het in de uitvoer.

Kernel Skelleton

Laten we door de meest eenvoudige kernel lopen die er is en enkele variaties daarop

__kernel void myKernel() {
}

Een kernel die vanuit uw hoofdcode kan worden gestart, wordt geïdentificeerd door het sleutelwoord __kernel. Een kernelfunctie kan alleen een retourtype ongeldig hebben.

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

}

Natuurlijk kun je meer functies maken die niet als kernels worden weergegeven. In dit geval kunt u de __kernel-modificator gewoon weglaten.

Een functie kan variabelen blootleggen zoals elke andere C / C ++ - functie. Het enige verschil is wanneer u naar het geheugen wilt verwijzen. Dit geldt voor alle verwijzingen, of het nu argumenten zijn of in code worden gebruikt.

float*  ptr;

is een pointer naar een geheugengebied waar alleen de uitvoerende thread toegang toe heeft. In feite is het hetzelfde als

__private float* ptr;

Er zijn vier verschillende modi voor geheugenregio's beschikbaar. In de kernel hoef je je er meestal geen zorgen over te maken, maar voor argumenten is dit essentieel.

• __global: deze modifier verwijst naar een pointer die in het globale geheugen wordt geplaatst
• __constant: verwijst naar een constante geheugenpointer
• __local: verwijst naar een gedeelde geheugenpointer
• __private: verwijst naar een lokale geheugenpointer

Bovendien kunnen we definiëren hoe we toegang tot het geheugen willen

• geen modificatie: lezen en schrijven
• __alleen lezen
• __write_only

Die vlaggen moeten overeenkomen met de manier waarop we de geheugenbuffer op onze host hebben toegewezen.

Kernel ID

Om goed met de gegevens te kunnen werken, moet elke thread zijn positie in de threadblock / globale threadpool kennen. Dit kan worden gearchiveerd met

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

Die twee functies retourneren de positie van de thread ten opzichte van het threadblock of alle threads.

get_working_dim();

Hiermee wordt het totale aantal dimensies opgehaald waarmee de kernel is gelanceerd.

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

Hiermee wordt het totale aantal threads in het threadblock of in totaal voor een bepaalde dimensie opgehaald.

Voorbehoud: zorg er altijd voor dat uw thread uw gegevensgrootte niet overschrijdt. Dit is zeer waarschijnlijk en moet altijd worden gecontroleerd.

Vectoren in OpenCL

Elk fundamenteel opencl-type heeft een vectorversie. U kunt het vectortype gebruiken door het aantal gewenste componenten na het type toe te voegen. Ondersteund aantal componenten zijn 2,3,4,8 en 16. OpenCL 1.0 biedt geen drie componenten.

U kunt elke vector op twee manieren initialiseren:

• Zorg voor een enkele scalaire
• Voldoe aan alle componenten

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

of

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

of elke andere combinatie van vectoren die voldoen aan het aantal componenten. Om toegang te krijgen tot de elementen van een vector kunt u verschillende methoden gebruiken. U kunt indexering gebruiken:

a[0] = 2;

of gebruik literals. Het voordeel van letterlijk is dat je ze kunt combineren zoals je wilt, doe dat in een moment. U hebt toegang tot alle vectorcomponenten met

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

u kunt ook meerdere componenten combineren tot een nieuwe vector

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

u kunt de volgorde van de componenten op elke gewenste manier wijzigen.

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

maar je kunt zoiets niet doen

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

Er zijn enkele handige steno voor toegang tot vectorcomponenten. De volgende steno's zijn alleen van toepassing op maat 2, 4, 8 en 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

Voor vectorgroottes 2,3 en 4 zijn er enkele extra steno

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

Gamma Correctie kernel

Laten we eens kijken naar een gamma-correctie kernel

__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);
    }
}

Laten we nu stap voor stap door die code lopen. De eerste regel maakt een variabele in het __constante geheugengebied van het type sampler_t. Deze sampler wordt gebruikt om de toegang tot onze afbeeldingsgegevens nader te specificeren. Raadpleeg de Khronos-documenten voor een volledige documentatie.

We hebben de input toegewezen als read_only en de output als write_only voordat we onze kernel aanroerden, dus we voegen die modificatoren hier toe.

image2d en image3d worden altijd toegewezen aan het wereldwijde geheugen, daarom kunnen we de __global-modifier hier weglaten. Onze gammawaarde bevindt zich in __constant geheugen, dus we specificeren dat ook.

Vervolgens krijgen we onze thread-ID die de pixel definieert die we naar gamma-correct gaan. We vragen ook naar de grootte om te controleren of onze thread geen toegang krijgt tot niet-toegewezen geheugen. Dit zal je kernel zeker laten crashen als je dat vergeet.

Nadat we ervoor hebben gezorgd dat we een legitieme thread zijn, lezen we onze pixel uit onze invoerafbeelding. We converteren het vervolgens naar zwevend om verlies van decimalen te voorkomen, voeren enkele berekeningen uit, zetten het terug en schrijven het in de uitvoer.