opencl
Kernel-Grundlagen

Einführung

Graustufen-Kernel

Wir können einen Kernel erstellen, um ein Graustufenbild zu erzeugen. Wir verwenden Bilddaten, die für jede Komponente mit Hilfe von uints definiert werden und mit der Reihenfolge 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);
    }
}

Gehen wir nun Schritt für Schritt durch diesen Code. Die erste Zeile erstellt eine Variable im __constant-Speicherbereich des Typs sampler_t. Dieser Sampler dient dazu, den Zugriff auf unsere Bilddaten weiter zu spezifizieren. Eine vollständige Dokumentation finden Sie in den Khronos-Dokumenten.

Wir haben die Eingabe als read_only und die Ausgabe als write_only zugewiesen, bevor wir unseren Kernel aufgerufen haben, sodass wir diese Modifikatoren hier hinzufügen.

image2d und image3d werden immer auf dem globalen Speicher zugewiesen, daher können wir hier den Modifizierer __global weglassen.

Dann erhalten wir unsere Thread-ID, die den Pixel definiert, den wir in Graustufen konvertieren wollen. Wir fragen auch nach der Größe, um sicherzustellen, dass unser Thread nicht auf nicht zugewiesenen Speicher zugreift. Dies wird Ihren Kernel definitiv zum Absturz bringen, wenn Sie das vergessen.

Nachdem wir sichergestellt haben, dass wir ein legitimer Thread sind, lesen wir unser Pixel aus unserem Eingabebild. Wir konvertieren es dann in Float, um den Verlust von Dezimalstellen zu vermeiden, führen Berechnungen durch, konvertieren es zurück und schreiben es in die Ausgabe.

Kernel Skelleton

Lass uns durch den einfachsten Kernel gehen, den es gibt, und einige Variationen davon

__kernel void myKernel() {
}

Ein Kernel, der vom Hauptcode aus gestartet werden kann, wird durch das Schlüsselwort __kernel identifiziert. Eine Kernel-Funktion kann nur den Rückgabetyp ungültig haben.

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

}

Natürlich können Sie weitere Funktionen erstellen, die nicht als Kernel verfügbar gemacht werden. In diesem Fall können Sie den Modifizierer __kernel einfach weglassen.

Eine Funktion kann Variablen wie jede andere C / C ++ - Funktion verfügbar machen. Der einzige Unterschied besteht darin, wann Sie auf Speicher zugreifen möchten. Dies gilt für alle Zeiger, unabhängig davon, ob sie Argumente sind oder im Code verwendet werden.

float*  ptr;

ist ein Zeiger auf einen Speicherbereich, zu dem nur der ausführende Thread Zugriff hat. In der Tat ist es das gleiche wie

__private float* ptr;

Es stehen vier verschiedene Speicherbereichsmodifizierer zur Verfügung. Innerhalb des Kernels müssen Sie sich normalerweise keine Sorgen machen, aber für Argumente ist dies unerlässlich.

• __global: Dieser Modifikator bezieht sich auf einen Zeiger, der sich im globalen Speicher befindet
• __constant: bezieht sich auf einen konstanten Speicherzeiger
• __local: bezieht sich auf einen Shared Memory-Zeiger
• __private: bezieht sich auf einen lokalen Speicherzeiger

Außerdem können wir festlegen, wie wir auf den Speicher zugreifen möchten

• Kein Modifikator: Lesen und Schreiben
• __schreibgeschützt
• __write_only

Diese Flags müssen mit der Art und Weise übereinstimmen, wie wir den Speicherpuffer auf unserem Host zugewiesen haben.

Kernel-ID

Um richtig mit den Daten arbeiten zu können, muss jeder Thread seine Position im Threadblock / globalen Threadpool kennen. Dies kann mit archiviert werden

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

Diese beiden Funktionen geben die Position des Threads relativ zum Threadblock oder zu allen Threads zurück.

get_working_dim();

Ruft die Gesamtzahl der Dimensionen ab, mit denen der Kernel gestartet wurde.

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

Ruft die Gesamtzahl der Threads im Threadblock oder insgesamt für eine bestimmte Dimension ab.

Vorbehalt: Stellen Sie immer sicher, dass Ihr Thread Ihre Datengröße nicht überschreitet. Dies ist sehr wahrscheinlich und sollte immer überprüft werden.

Vektoren in OpenCL

Jeder grundlegende opencl-Typ hat eine Vektorversion. Sie können den Vektortyp verwenden, indem Sie die Anzahl der gewünschten Komponenten nach dem Typ anhängen. Unterstützte Anzahl von Komponenten sind 2,3,4,8 und 16. OpenCL 1.0 bietet keine drei Komponenten.

Sie können jeden Vektor auf zwei Arten initialisieren:

• Geben Sie einen einzelnen Skalar an
• Erfüllen Sie alle Komponenten

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

oder

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

oder eine beliebige andere Kombination von Vektoren, die die Anzahl der Komponenten erfüllen. Um auf die Elemente eines Vektors zuzugreifen, können Sie verschiedene Methoden verwenden. Sie können entweder die Indizierung verwenden:

a[0] = 2;

oder verwenden Sie Literale. Der Vorteil von Literalen besteht darin, dass Sie sie beliebig kombinieren können. Mit können Sie auf alle Vektorkomponenten zugreifen

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

Sie können auch mehrere Komponenten zu einem neuen Vektor kombinieren

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

Sie können die Reihenfolge der Komponenten beliebig ändern.

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

aber so etwas kann man nicht machen

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

Es gibt einige praktische Abkürzungen für den Zugriff auf Vektorkomponenten. Die folgenden Abkürzungen gelten nur für die Größen 2, 4, 8 und 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 die Vektorgrößen 2, 3 und 4 gibt es zusätzliche Abkürzungen

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

Gamma-Korrekturkern

Schauen wir uns einen Gamma-Korrekturkern an

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

Gehen wir nun Schritt für Schritt durch diesen Code. Die erste Zeile erstellt eine Variable im __constant-Speicherbereich des Typs sampler_t. Dieser Sampler dient dazu, den Zugriff auf unsere Bilddaten weiter zu spezifizieren. Eine vollständige Dokumentation finden Sie in den Khronos-Dokumenten.

Wir haben die Eingabe als read_only und die Ausgabe als write_only zugewiesen, bevor wir unseren Kernel aufgerufen haben, sodass wir diese Modifikatoren hier hinzufügen.

image2d und image3d werden immer auf dem globalen Speicher zugewiesen, daher können wir hier den Modifizierer __global weglassen. Unser Gamma-Wert befindet sich im __constant-Speicher. Daher geben wir dies auch an.

Dann erhalten wir unsere Thread-ID, die den Pixel definiert, den wir für die Gamma-Korrektur verwenden. Wir fragen auch nach der Größe, um sicherzustellen, dass unser Thread nicht auf nicht zugewiesenen Speicher zugreift. Dies wird Ihren Kernel definitiv zum Absturz bringen, wenn Sie das vergessen.

Nachdem wir sichergestellt haben, dass wir ein legitimer Thread sind, lesen wir unser Pixel aus unserem Eingabebild. Wir konvertieren es dann in Float, um den Verlust von Dezimalstellen zu vermeiden, führen Berechnungen durch, konvertieren es zurück und schreiben es in die Ausgabe.