opencl
Nozioni di base sul kernel

introduzione

Kernel in scala di grigi

Consente di creare un kernel per generare un'immagine in scala di grigio. Useremo i dati dell'immagine che sono definiti usando gli uints per ciascun componente e con l'ordine 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);
    }
}

Ora passiamo attraverso quel codice passo dopo passo. La prima riga crea una variabile nella regione di memoria __constant di tipo sampler_t. Questo campionatore viene utilizzato per specificare ulteriormente l'accesso ai nostri dati di immagine. Si prega di fare riferimento ai documenti di Khronos per una documentazione completa.

Abbiamo assegnato l'input come read_only e l'output come write_only prima di chiamare il nostro kernel, quindi aggiungiamo questi modificatori qui.

image2d e image3d sono sempre allocati nella memoria globale, quindi qui possiamo omettere il modificatore __global.

Quindi otteniamo il nostro ID thread che definisce il pixel che convertiamo in scala di grigi. Esaminiamo anche la dimensione per assicurarci che il nostro thread non stia accedendo alla memoria non allocata. Questo sicuramente arresterà il tuo kernel se lo dimentichi.

Dopo esserci assicurati di essere un thread legittimo, leggiamo il nostro pixel fuori dalla nostra immagine di input. Quindi lo convertiamo in float per evitare la perdita di posizioni decimali, fare alcuni calcoli, riconvertirlo e scriverlo nell'output.

Kernel Skelleton

Passiamo attraverso il kernel più semplice che ci sia e alcune varianti di esso

__kernel void myKernel() {
}

Un kernel che può essere avviato dal tuo codice principale è identificato dalla parola chiave __kernel. Una funzione del kernel può avere solo il tipo di ritorno vuoto.

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

}

Naturalmente puoi creare più funzioni che non sono esposte come kernel. In questo caso puoi semplicemente omettere il modificatore __kernel.

Una funzione può esporre variabili come qualsiasi altra funzione C / C ++. L'unica differenza è quando vuoi fare riferimento alla memoria. Questo vale per tutti i puntatori sia che siano argomenti o usati nel codice.

float*  ptr;

è un puntatore a un'area di memoria a cui ha accesso solo il thread in esecuzione. In effetti è lo stesso di

__private float* ptr;

Sono disponibili quattro diversi modificatori della regione di memoria. All'interno del kernel di solito non devi preoccuparti di questo, ma per argomenti questo è essenziale.

• __global: questo modificatore si riferisce a un puntatore posizionato nella memoria globale
• __constant: si riferisce a un puntatore di memoria costante
• __local: si riferisce a un puntatore di memoria condivisa
• __private: fa riferimento a un puntatore di memoria locale

Inoltre possiamo definire come vogliamo accedere alla memoria

• nessun modificatore: leggi e scrivi
• __sola lettura
• __write_only

Questi flag devono corrispondere al modo in cui abbiamo allocato il buffer di memoria sul nostro host.

Kernel ID

Per lavorare correttamente con i dati, ogni thread deve conoscere la sua posizione nel threadblock / pool di thread globale. Questo può essere raggiunto con

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

Queste due funzioni restituiscono la posizione del thread relativa al threadblock o a tutti i thread.

get_working_dim();

Ottiene il numero totale di dimensioni con cui è stato avviato il kernel.

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

Ottiene il numero totale di thread nel threadblock o in totale per una determinata dimensione.

Avvertenza: accertati sempre che il tuo thread non superi la dimensione dei tuoi dati. È molto probabile che ciò accada e dovrebbe sempre essere controllato.

Vettori in OpenCL

Ogni tipo di opencl fondamentale ha una versione vettoriale. È possibile utilizzare il tipo di vettore aggiungendo il numero di componenti desiderati dopo il tipo. Il numero di componenti supportato è 2,3,4,8 e 16. OpenCL 1.0 non offre tre componenti.

Puoi inizializzare qualsiasi vettore usando due modi:

• Fornire un singolo scalare
• Soddisfa tutti i componenti

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

o

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

o qualsiasi altra combinazione di vettori che soddisfano il numero di componenti. Per accedere agli elementi di un vettore è possibile utilizzare diversi metodi. Puoi utilizzare l'indicizzazione:

a[0] = 2;

o usare letterali. Il vantaggio dei letterali è che puoi combinarli come vuoi, e farlo in un attimo. È possibile accedere a tutti i componenti vettoriali con

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

puoi anche combinare più componenti in un nuovo vettore

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

puoi cambiare l'ordine dei componenti nel modo che preferisci.

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

ma non puoi fare qualcosa del genere

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

Esistono alcune pratiche abbreviazioni per l'accesso ai componenti vettoriali. Le seguenti stenografie si applicano solo alle taglie 2, 4, 8 e 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

Per le dimensioni dei vettori 2,3 e 4 esistono alcune abbreviazioni aggiuntive

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

Gamma Correzione del kernel

Vediamo un kernel di correzione della 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);
    }
}

Ora passiamo attraverso quel codice passo dopo passo. La prima riga crea una variabile nella regione di memoria __constant di tipo sampler_t. Questo campionatore viene utilizzato per specificare ulteriormente l'accesso ai nostri dati di immagine. Si prega di fare riferimento ai documenti di Khronos per una documentazione completa.

Abbiamo assegnato l'input come read_only e l'output come write_only prima di chiamare il nostro kernel, quindi aggiungiamo questi modificatori qui.

image2d e image3d sono sempre allocati nella memoria globale, quindi qui possiamo omettere il modificatore __global. Il nostro valore gamma si trova nella memoria __ costante, quindi lo specificiamo anche.

Quindi otteniamo il nostro ID thread che definisce il pixel che andremo a correggere gamma. Esaminiamo anche la dimensione per assicurarci che il nostro thread non stia accedendo alla memoria non allocata. Questo sicuramente arresterà il tuo kernel se lo dimentichi.

Dopo esserci assicurati di essere un thread legittimo, leggiamo il nostro pixel fuori dalla nostra immagine di input. Quindi lo convertiamo in float per evitare la perdita di posizioni decimali, fare alcuni calcoli, riconvertirlo e scriverlo nell'output.