====== einfache Vektoraddition ======

Der Sinn des Programms ist es, zu zeigen, wie OpenCL grundsätzlich funktioniert, was beachtet werden muss. Weiter unten wird dann alles erklärt. 

===== Quellcode =====

<code cpp host.c>
#include <stdio.h>
#include <stdlib.h>
 
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
 
#define MAX_SOURCE_SIZE (0x100000)
 
int main(void) {
    // Create the two input vectors
    int i;
    const int LIST_SIZE = 1024;
    int *A = (int*)malloc(sizeof(int)*LIST_SIZE);
    int *B = (int*)malloc(sizeof(int)*LIST_SIZE);
    for(i = 0; i < LIST_SIZE; i++) {
        A[i] = i;
        B[i] = LIST_SIZE - i;
    }
 
    // Load the kernel source code into the array source_str
    FILE *fp;
    char *source_str;
    size_t source_size;
 
    fp = fopen("vector_add_kernel.cl", "r");
    if (!fp) {
        fprintf(stderr, "Failed to load kernel.\n");
        exit(1);
    }
    source_str = (char*)malloc(MAX_SOURCE_SIZE);
    source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
    fclose( fp );
 
    // Get platform and device information
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;   
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_DEFAULT, 1, 
            &device_id, &ret_num_devices);
 
    // Create an OpenCL context
    cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret);
 
    // Create a command queue
    cl_command_queue command_queue = clCreateCommandQueue(context, 
            device_id, 0, &ret);
 
    // Create memory buffers on the device for each vector 
    cl_mem a_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, 
            LIST_SIZE * sizeof(int), NULL, &ret);
    cl_mem b_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY,
            LIST_SIZE * sizeof(int), NULL, &ret);
    cl_mem c_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
            LIST_SIZE * sizeof(int), NULL, &ret);
 
    // Copy the lists A and B to their respective memory buffers
    ret = clEnqueueWriteBuffer(command_queue, a_mem_obj, CL_TRUE, 0,
            LIST_SIZE * sizeof(int), A, 0, NULL, NULL);
    ret = clEnqueueWriteBuffer(command_queue, b_mem_obj, CL_TRUE, 0, 
            LIST_SIZE * sizeof(int), B, 0, NULL, NULL);
 
    // Create a program from the kernel source
    cl_program program = clCreateProgramWithSource(context, 1, 
            (const char **)&source_str, (const size_t *)&source_size, &ret);
 
    // Build the program
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
 
    // Create the OpenCL kernel
    cl_kernel kernel = clCreateKernel(program, "vector_add", &ret);
 
    // Set the arguments of the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_mem_obj);
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_mem_obj);
 
    // Execute the OpenCL kernel on the list
    size_t global_item_size = LIST_SIZE; // Process the entire lists
    size_t local_item_size = 64; // Divide work items into groups of 64
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
            &global_item_size, &local_item_size, 0, NULL, NULL);
 
    // Read the memory buffer C on the device to the local variable C
    int *C = (int*)malloc(sizeof(int)*LIST_SIZE);
    ret = clEnqueueReadBuffer(command_queue, c_mem_obj, CL_TRUE, 0, 
            LIST_SIZE * sizeof(int), C, 0, NULL, NULL);
 
    // Display the result to the screen
    for(i = 0; i < LIST_SIZE; i++)
        printf("%d + %d = %d\n", A[i], B[i], C[i]);
 
    // Clean up
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(a_mem_obj);
    ret = clReleaseMemObject(b_mem_obj);
    ret = clReleaseMemObject(c_mem_obj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
    free(A);
    free(B);
    free(C);
    return 0;
}
</code>

<code c vector_add_kernel.cl>
__kernel void vector_add(__global const float *a,
                                                __global const float *b,
                                                __global float *result)
{
    int gid = get_global_id(0);
 
    result[gid] = a[gid] + b[gid];
}
</code>

<code bash>
gcc host.c -lOpenCL
</code>

===== Erläuterung zum Quellcode =====

Das C-Programm initialisiert die OpenCL-Umgebung, liest den Kernel ein, holt sich das Device, erstellt einen Kontext und eine CommandQueue, die Daten auf die Grafikkarte (das OpenCL-Gerät) kopiert, dort werden die Daten im Kernel verarbeitet, und nach Beendigung des Kernels im Hostprogramm zurückkopiert.

==== Devices und Plattformen ====


<code c>
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;   
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_DEFAULT, 1, 
            &device_id, &ret_num_devices);
</code>

Zunächst werden Variablen für Plattform ID und Device ID angelegt, sowie Variablen um die Anzahl der Devices und Plattformen zu halten. In OpenCL ist es möglich mehrere Plattformen in einem Computer zu haben, bspw. kann eine AMD Grafikkarte und eine Intel CPU Verwendung finden. AMD und Intel liefern unterschiedliche OpenCL-Umgebungen, und man kann innerhalb des Programms zwischen den vorhandenen Plattformen wählen, oder sich irgendeine nehmen.

''clGetPlatformIDs'' lässt sich vom OpenCL-Framework eine Plattform ID geben, und ''clGetDeviceIDs'' lässt sich eine Device ID geben. Genaue Beschreibungen der Funktionen finden Sie in unserer Referenz oder unter [1]. Es gibt die Möglichkeit ''clGetDeviceIDs'' mitzuteilen, nur bestimmte Geräte haben zu wollen. Zum Einen muss man sich auf eine Plattform festlegen, zum anderen kann man angeben, ob CPUs, GPUs, andere Beschleuniger gewählt werden sollen oder das egal ist (DEFAULT => CPU oder GPU).

==== Kontext und Commandqueue ====

<code c>
    // Create an OpenCL context
    cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret);
 
    // Create a command queue
    cl_command_queue command_queue = clCreateCommandQueue(context, 
            device_id, 0, &ret);
</code>

Als nächstes erzeugen wir uns aus der Plattform ID und der DeviceID einen Kontext und eine Commandqueue. Der Kontext lässt uns auf das Device zugreifen, die Commandqueue ist eine Befehlswarteschlange. Im Laufe des Programms werden wir Befehle an das OpenCL-Device schicken (und damit Operationen in die Warteschlange einreihen, die vom Gerät ausgeführt werden sollen). 

==== Speicherobjekte ====

<code c>
    // Create memory buffers on the device for each vector 
    cl_mem a_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, 
            LIST_SIZE * sizeof(int), NULL, &ret);
    cl_mem b_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY,
            LIST_SIZE * sizeof(int), NULL, &ret);
    cl_mem c_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
            LIST_SIZE * sizeof(int), NULL, &ret);
</code>

Nun werden Speicherobjekte angelegt - das kann man ähnlich auffassen, wie Variablen (bzw. Felder / Zeiger) in einem normalen C-Programm. Wir legen die Speicherobjekte an, was veranlasst, dass die OpenCL-Umgebung den Speicher auf dem Device reserviert, oder abbricht, wenn er nicht zur Verfügung steht. Die Hostvariable ''a_mem_obj'' des Typs ''cl_mem'' ist ein Identifikator, mit dem wir dem OpenCL-Framework mitteilen können, welches Speicherobjekt angesprochen werden soll. Nach der Benutzung von Speicherobjekten ist es erforderlich diese auch wieder zu löschen (vgl. malloc und free). Wenn man einem Compute Kernel Felder als Parameter übergeben möchte, ist es erforderlich Speicherobjekte anzulegen. 

<code c>
    // Copy the lists A and B to their respective memory buffers
    ret = clEnqueueWriteBuffer(command_queue, a_mem_obj, CL_TRUE, 0,
            LIST_SIZE * sizeof(int), A, 0, NULL, NULL);
    ret = clEnqueueWriteBuffer(command_queue, b_mem_obj, CL_TRUE, 0, 
            LIST_SIZE * sizeof(int), B, 0, NULL, NULL);
</code>

''clEnqueueWriteBuffer'' fügt eine Schreiboperation auf ein Speicherobjekt in die Befehlswarteschlange ein. Dh. in der Befehlswarteschlange steht nach der Ausführung der Funktion eine Schreiboperation, die später vom CL-Device ausgeführt wird. Dabei muss angegeben werden, welches Speicherobjekt beschrieben werden soll, wie viele Bytes geschrieben werden sollen und wo die Daten herkommen, sowie die Command-Queue Variable. 

''clEnqueueWriteBuffer'' führt den Schreibbefehl nicht sofort aus, sondern fügt ihn nur in die Befehlwarteschlange ein. 

==== Kernelerzeugung ====

<code c>
    // Create a program from the kernel source
    cl_program program = clCreateProgramWithSource(context, 1, 
            (const char **)&source_str, (const size_t *)&source_size, &ret);
 
    // Build the program
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
 
    // Create the OpenCL kernel
    cl_kernel kernel = clCreateKernel(program, "vector_add", &ret);
</code>

Wir haben weiter oben einen String aus einer Datei ausgelesen. Dieser String muss von einem OpenCL-Compiler übersetzt werden, das geschrieht mit ''clCreateProgramWithSource'' und ''clBuildProgram''. In diesem Code (weiter unten mehr dazu) gibt es eine Funktion, die mit kernel gekennzeichnet ist. Diese Funktion ist der eigentliche Kernel, der vom Hostprogramm aus gerufen werden kann. Dieser Kernel wird mittels ''clCreateKernel'' erzeugt.

==== Parameter des Kernels ====

<code c>
    // Set the arguments of the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_mem_obj);
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_mem_obj);
</code>

Die Kernelfunktion besitzt drei Parameter. Diese Parameter müssen im Hostprogramm festgelegt werden. Das geschieht mittels ''clSetKernelArg''. Dabei wird das Kernel-Objekt angegeben, die Parameternummer, und die Größe des eigentlichen Parameters, sowie der Parameter selbst. Wichtig dabei ist, dass, wenn ein Speicherobjekt als Parameter verwendet wird, ''sizeof(cl_mem)'' als Größe angegeben wird. Soll bspw. ein Int-Wert übergeben werden, kann dort sizeof(int) stehen. 

==== Ausführung ====

<code c>
    // Execute the OpenCL kernel on the list
    size_t global_item_size = LIST_SIZE; // Process the entire lists
    size_t local_item_size = 64; // Divide work items into groups of 64
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
            &global_item_size, &local_item_size, 0, NULL, NULL);
</code>

''clEnqueueNDRangeKernel'' reiht die Ausführung des Kernels in die Befehlswarteschlange des CL-Devices ein. Für die Ausführung eines Kernel müssen die Ausmaße der Berechnung bekannt sein (Parallelität). Wir haben hier eine LIST_SIZE große Datenstruktur, die parallel bearbeitet werden kann. Deshalb verwenden wir LIST_SIZE viele Working Groups (Ausführungseinheiten). local_item_size gibt dann noch an, wie groß die Working Groups sein sollen. Es ist auch möglich mehrere Dimensionen von Ausführungseinheiten anzulegen, das erleichtert bspw. die Programmierung für die Bearbeitung eines Bildes.

==== Daten zurücklesen ====

<code c>
    // Read the memory buffer C on the device to the local variable C
    int *C = (int*)malloc(sizeof(int)*LIST_SIZE);
    ret = clEnqueueReadBuffer(command_queue, c_mem_obj, CL_TRUE, 0, 
            LIST_SIZE * sizeof(int), C, 0, NULL, NULL);
</code>

''clEnqueueReadBuffer'' funktioniert analog zu ''clEnqueueWriteBuffer''. Hier wird eine Leseoperation in die Befehlswarteschlange eingefügt, die den Wert des Speicherobjekts c_mem_obj ausliest und an die Speicheradresse C des Host-Speichers schreibt. 

==== Aufräumen ====

<code c>
    // Clean up
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(a_mem_obj);
    ret = clReleaseMemObject(b_mem_obj);
    ret = clReleaseMemObject(c_mem_obj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
    free(A);
    free(B);
    free(C);
    return 0;
</code>

Hier wird aufgeräumt. ClFlush und clFinish haben zum Ergebnis, dass die Befehlswarteschlange abgearbeitet wird (also zuende abgearbeitet wird, clFlush wartet nicht darauf, clFinish wartet auf die Abarbeitung der CommandQueue.

clReleaseX leert den Speicher, den ein Objekt eingenommen hat, bspw. wird bei clReleaseMemObject der Speicher auf dem Gerät freigegeben, der vom Speicherobjekt eingenommen wurde. Nicht vergessen den Kontext und die CommandQueue wieder freizugeben. 

===== Kernel =====

Der Kernel ist der eigentlich auf dem CL-Device ausgeführte Code. Eine Kerneldatei endet auf .cl und sieht erstmal aus wie normaler C-Code. Folgender Kernel: 

<code c vector_add_kernel.cl>
__kernel void vector_add  (__global const float *a,
                           __global const float *b,
                           __global float *result )
{
    int gid = get_global_id(0);
 
    result[gid] = a[gid] + b[gid];
}
</code>

Zunächst steht der Funktionskopf, wie bei einer normalen C-Funktion. Das beginnende kernel hat zur Folge, dass die Funktion als Kernel vom Hauptprogramm aus aufrufbar ist, sie dient quasi als Einsprungpunkt für die Berechhnung auf dem CL-Device. Es ist möglich weitere Funktionen zu haben, die auf dem CL-Device ausgeführt werden können, die diese Markierung mittels kernel nicht besitzen. 

Es folgt ein Rückgabetyp und Funktionsname, sowie die Parameterliste. Es ist auffällig, dass die Parameter einerseits ''global'' andererseits (die ersten beiden) als const definiert sind. ''const'' verhindert, dass die Variable verändert werden kann, so können Fehler beim Programmieren vermieden werden, wenn der Programmierer doch versucht die Variable zu ändern. ''global'' bedeutet, dass die Variablen im (langsameren) globalen Speicher liegen. Es gibt noch ''local'', dann liegen die Parameter im lokalen Speicher, der schneller, aber kleiner ist. Außerdem steht der nur den lokalen Recheneinheiten zur Verfügung.

Der Kernel besteht nur aus zwei Anweisungen. Zunächst wird die globale ID des Working-Items (Ausführungseinheit) geholt. In Abhängigkeit dieses Wertes kann bestimmt werden, welche Stelle des Vektors von dem Working-Item bearbeitet werden soll. Kernelfunktionen werden also aus der Sicht irgendeines einzigen Working-Item geschrieben, und lässt sich anhand der globalen ID problemlos von Cl auf alle Berechnungseinheiten übertragen. 

===== Literatur =====

[1] [[https://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/]]