opencl
Kernel Basics

Introduktion

Gråskala kärna

Låter bygga en kärna för att generera en gråskalabild. Vi kommer att använda bilddata som definieras med hjälp av uints för varje komponent och med ordning 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);
    }
}

Låt oss nu gå igenom koden steg för steg. Den första raden skapar en variabel i __konstantminnesregionen av typ sampler_t. Denna samplare används för att ytterligare specificera åtkomsten till våra bilddata. Se Khronos-dokumenten för en fullständig dokumentation.

Vi tilldelade ingången som read_only och outputen som write_only innan vi kallade vår kärna så vi lägger till dessa modifierare här.

image2d och image3d tilldelas alltid i det globala minnet, därför kan vi utelämna __global modifieraren här.

Sedan får vi vårt tråd-id som definierar pixeln vi ska konvertera till gråskala. Vi frågar också storleken för att se till att vår tråd inte får åtkomst till oallokerat minne. Detta kommer definitivt att krascha din kärna om du glömmer det.

När vi såg till att vi är en legitim tråd, läste vi vår pixel ur vår inmatningsbild. Vi konverterar sedan det till flyta för att undvika förlust av decimaler, gör några beräkningar, konverterar tillbaka och skriver det till utgången.

Kernel Skelleton

Låter oss gå igenom den mest enkla kärnan som finns och vissa variationer av den

__kernel void myKernel() {
}

En kärna som kan startas från din huvudkod identifieras med nyckelordet __kernel. En kärnfunktion kan endast ha returtyp.

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

}

Naturligtvis kan du skapa fler funktioner som inte exponeras som kärnor. I det här fallet kan du bara utelämna __kärnmodifieraren.

En funktion kan exponera variabler som alla andra C / C ++ -funktioner skulle göra. Den enda skillnaden är när du vill referensminne. Detta gäller alla pekare oavsett om de är argument eller används i kod.

float*  ptr;

är en pekare till ett minnesområde endast den exekverande tråden har åtkomst till. I själva verket är det samma som

__private float* ptr;

Det finns fyra olika minnesregionmodifierare tillgängliga. Inuti kärnan behöver du vanligtvis inte oroa dig för det, men för argument är detta viktigt.

• __global: Den här modifieraren hänvisar till en pekare som placeras i det globala minnet
• __konstant: avser en konstant minnespekare
• __local: hänvisar till en delad minnespekare
• __privat: hänvisar till en lokal minnespekare

Dessutom kan vi definiera hur vi vill komma åt minnet

• ingen modifierare: läs och skriv
• __read_only
• __write_only

Dessa flaggor måste matcha det sätt som vi tilldelade minnesbufferten tillbaka till vår värd.

Kärn-ID

För att fungera korrekt med uppgifterna måste varje tråd veta sin position i trådblock / global trådpool. Detta kan arkiveras med

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

Dessa två funktioner returnerar trådens position relativt trådblocket eller alla trådar.

get_working_dim();

Hämtar det totala antalet dimensioner som kärnan lanserades med.

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

Hämtar det totala antalet trådar i trådblocket eller totalt för en viss dimension.

Varning: Se alltid till att din tråd inte överskrider din datastorlek. Det händer mycket troligt och bör alltid kontrolleras.

Vektorer i OpenCL

Varje grundläggande opencl-typ har en vektorversion. Du kan använda vektortypen genom att lägga till antalet önskade komponenter efter typen. Antalet komponenter som stöds är 2,3,4,8 och 16. OpenCL 1.0 erbjuder inte tre komponenter.

Du kan initialisera vilken vektor som helst på två sätt:

• Ge en enda skalär
• Tillfredsställa alla komponenter

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

eller

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

eller någon annan kombination av vektorer som uppfyller antalet komponenter. För att få tillgång till elementen i en vektor kan du använda olika metoder. Du kan antingen använda indexering:

a[0] = 2;

eller använd bokstäver. Fördelen med bokstäver är att du kan kombinera dem som du vill, och gör det på ett ögonblick. Du kan komma åt alla vektorkomponenter med

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

du kan också kombinera flera komponenter till en ny vektor

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

du kan ändra ordning på komponenterna på vilket sätt du vill.

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

men du kan inte göra något liknande

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

Det finns några bekväma korthår för att få tillgång till vektorkomponenter. Följande korthår gäller endast storlek 2, 4, 8 och 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

För vektorstorlekar 2,3 och 4 finns det några ytterligare korthår

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

Gamma-korrigeringskärna

Låt oss titta på en gammakorrektionskärna

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

Låt oss nu gå igenom koden steg för steg. Den första raden skapar en variabel i __konstantminnesregionen av typ sampler_t. Denna samplare används för att ytterligare specificera åtkomsten till våra bilddata. Se Khronos-dokumenten för en fullständig dokumentation.

Vi tilldelade ingången som read_only och outputen som write_only innan vi kallade vår kärna så vi lägger till dessa modifierare här.

image2d och image3d tilldelas alltid i det globala minnet, därför kan vi utelämna __global modifieraren här. Vårt gammavärde finns i __konstant minne, så vi anger det också.

Sedan får vi vårt tråd-id som definierar pixeln vi ska gå korrekt. Vi frågar också storleken för att se till att vår tråd inte får åtkomst till oallokerat minne. Detta kommer definitivt att krascha din kärna om du glömmer det.

När vi såg till att vi är en legitim tråd, läste vi vår pixel ur vår inmatningsbild. Vi konverterar sedan det till flyta för att undvika förlust av decimaler, gör några beräkningar, konverterar tillbaka och skriver det till utgången.