Spieleentwicklung mit OpenCL

Asteroidenhagel selbstgemacht
Kommentare

In diesem Artikel befassen wir uns mit den Grundlagen der OpenCL-Programmierung und zeigen die verschiedenen Einsatzmöglichkeiten im Rahmen der Spieleentwicklung.

In diesem Artikel befassen wir uns mit den Grundlagen der OpenCL-Programmierung und zeigen die verschiedenen Einsatzmöglichkeiten im Rahmen der Spieleentwicklung.

Quellcode auf CD
Zur Veranschaulichung der in diesem Artikel behandelten Themen, finden Sie auf der Leser-CD verschiedene Programmbeispiele und die aktuelle Version des den Beispielen zugrunde liegenden OpenGL/OpenCL-Frameworks.



OpenCL (Open Computing Language) könnte nun einen möglichen Ausweg aus diesem Dilemma bieten, denn es handelt sich hierbei um die zurzeit einzig verfügbare, gleichermaßen von Intel, AMD und NVIDIA unterstützte Technologie, mit deren Hilfe sich komplexe Berechnungen wahlweise auf der CPU oder der GPU durchführen lassen. Entwickelt wurde die OpenCL-Schnittstelle samt der zugehörigen Sprache OpenCL C von der Firma Apple in Zusammenarbeit mit IBM, AMD, Intel und NVIDIA. Am 8. Dezember 2008 erfolgte die Bekanntgabe der OpenCL-Spezifikation 1.0 durch die Khronos-Gruppe.

Die aktuelle Spezifikation 1.2, die nach wie vor vollständige Abwärtskompatibilität gewährleistet, wurde am 16. November 2011 veröffentlicht. Der grundlegende Aufbau einer OpenCL-Anwendung ist in Abbildung 1 skizziert. Die heutige und zukünftige Bedeutung der OpenCL-Technologie bemisst sich nicht zuletzt daran, dass sowohl Intel als auch AMD und NVIDIA mittlerweile jeweils ein eigenes Software Development Kit (SDK) zum kostenlosen Download anbieten. Die Programmbeispiele auf der Leser-CD wurden beispielsweise mit dem AMD APP SDK entwickelt, da dieses sowohl die Programmierung von OpenCL-Anwendungen für AMD-GPUs als auch für AMD und Intel CPUs ermöglicht. Hinzu kommt, dass die von AMD zur Verfügung gestellte OpenCL Runtime im Vergleich zur Intel OpenCL Runtime deutlich performantere Berechnungen auf der CPU ermöglicht, zumindest was die Programmbeispiele betrifft.

Abb. 1: Aufbau einer OpenCL-Anwendung

Einsatzmöglichkeiten in der Spieleentwicklung

Zunächst wurden die Einsatzmöglichkeiten von CUDA und OpenCL vorwiegend im wissenschaftlichen Bereich gesehen, denn durch die Ausnutzung der Leistungsreserven moderner Grafikkarten lassen sich die für komplexe Simulationen erforderlichen Rechenzeiten (Wetter, Astronomie etc.) erheblich verkürzen (GPGPU: General Purpose Computation On Graphics Processing Unit). Auch im Bereich Multimedia gibt es immer neue Anwendungsmöglichkeiten, beispielsweise bei der Bearbeitung von Bildern und Filmen. Für die Spieleentwicklung ist der Einsatz von OpenCL insbesondere aus den folgenden drei Gründen sehr interessant:

  • Die optimale Ausnutzung der Leistungsreserven moderner Mehrkernprozessoren wird erleichtert.
  • Auf die Entwicklung eines eigenen Frameworks für die Verteilung der anstehenden Berechnungen auf die zur Verfügung stehenden CPU-Kerne könnte zukünftig verzichtet werden.
  • Berechnungen lassen sich flexibel zwischen CPU und GPU aufteilen.

Die konkreten Einsatzmöglichkeiten sind von der Art und dem Umfang des betreffenden Spiels abhängig. Betrachten wir hierzu einige Beispiele:

  • Die Entwicklung einer Physik-Engine, die dank OpenCL alle zur Verfügung stehenden CPU-Kerne optimal ausnutzt und zudem wahlweise Berechnungen auf der CPU und der GPU durchführen kann (z. B. Fluidpartikelsimulation/Wasser auf der GPU, spielrelevante Physik auf der CPU).
  • Kollisionsprävention (1): Eine Flotte von Raumschiffen muss möglichst unbeschadet ein Asteroidenfeld durchqueren. Um nicht mit anderen Schiffen oder Asteroiden zu kollidieren, müssen die einzelnen Raumschiffe von Zeit zu Zeit ihre Flugrichtung ändern.
  • Kollisionsprävention (2): In einem urbanen Szenario bevölkern Fußgänger die Gehwege und bewegen sich dabei mit unterschiedlichen Geschwindigkeiten in unterschiedliche Richtungen. Von Zeit zu Zeit müssen die Fußgänger die Laufrichtung ändern, um nicht mit anderen Passanten zusammenzustoßen.
  • In einer Weltraum-Wirtschafts-Simulation sollen für alle bewohnten Sonnensysteme die Marktpreise der Handelswaren auf Basis von Verfügbarkeit und Nachfrage dynamisch festgelegt werden. Eine erhöhte Nachfrage führt in den betreffenden Systemen zu einem Preisanstieg; in anderen Systemen führt ein Überangebot hingegen zu fallenden Preisen.

Wie sinnvoll der Einsatz von OpenCL für eine bestimmte Aufgabe ist, hängt entscheidend davon ab, ob und wie sich besagtes Problem parallelisieren lässt. Beim so genannten „aufgabenparallelen Multithreading“ wird eine Aufgabe als Ganzes in einem separaten Thread abgearbeitet (Beispiel: Streaming von Audio-, Geometrie- und Texturdaten). Beim „datenparallelen Multithreading“ wird eine Aufgabe hingegen in unabhängige Teilaufgaben zerlegt, die dann parallel zueinander in mehreren Threads bearbeitet werden. Auf die bekannten Werkzeuge zur Synchronisierung sollte hierbei wenn möglich verzichtet werden, denn Threads, die lange aufeinander warten müssen oder sich häufig gegenseitig blockieren, verringern den Vorteil der Parallelisierung.

Für alle Aufgaben, die sich datenparallel abarbeiten lassen, so auch für die zuvor besprochenen Beispiele (Teilaufgaben: Bestimme die Flugrichtung von Raumschiff i; Bestimme die Laufrichtung von Fußgänger i; Bestimme die Marktpreise für Sonnensystem i etc.), bietet sich der Einsatz von OpenCL an. Als Programmierer müssen wir uns nicht darum kümmern, wie viele Prozessor-Threads oder Shader-Einheiten auf den jeweiligen Rechnern zur Verfügung stehen, denn für die Aufteilung der Berechnungen ist die OpenCL Runtime verantwortlich. Wir können uns also uneingeschränkt mit der Parallelisierung der eigentlichen Aufgabe befassen.

  • Die Entwicklung einer Physik-Engine, die dank OpenCL alle zur Verfügung stehenden CPU-Kerne optimal ausnutzt und zudem wahlweise Berechnungen auf der CPU und der GPU durchführen kann (z. B. Fluidpartikelsimulation/Wasser auf der GPU, spielrelevante Physik auf der CPU).
  • Kollisionsprävention (1): Eine Flotte von Raumschiffen muss möglichst unbeschadet ein Asteroidenfeld durchqueren. Um nicht mit anderen Schiffen oder Asteroiden zu kollidieren, müssen die einzelnen Raumschiffe von Zeit zu Zeit ihre Flugrichtung ändern.
  • Kollisionsprävention (2): In einem urbanen Szenario bevölkern Fußgänger die Gehwege und bewegen sich dabei mit unterschiedlichen Geschwindigkeiten in unterschiedliche Richtungen. Von Zeit zu Zeit müssen die Fußgänger die Laufrichtung ändern, um nicht mit anderen Passanten zusammenzustoßen.
  • In einer Weltraum-Wirtschafts-Simulation sollen für alle bewohnten Sonnensysteme die Marktpreise der Handelswaren auf Basis von Verfügbarkeit und Nachfrage dynamisch festgelegt werden. Eine erhöhte Nachfrage führt in den betreffenden Systemen zu einem Preisanstieg; in anderen Systemen führt ein Überangebot hingegen zu fallenden Preisen.

Wie sinnvoll der Einsatz von OpenCL für eine bestimmte Aufgabe ist, hängt entscheidend davon ab, ob und wie sich besagtes Problem parallelisieren lässt. Beim so genannten „aufgabenparallelen Multithreading“ wird eine Aufgabe als Ganzes in einem separaten Thread abgearbeitet (Beispiel: Streaming von Audio-, Geometrie- und Texturdaten). Beim „datenparallelen Multithreading“ wird eine Aufgabe hingegen in unabhängige Teilaufgaben zerlegt, die dann parallel zueinander in mehreren Threads bearbeitet werden. Auf die bekannten Werkzeuge zur Synchronisierung sollte hierbei wenn möglich verzichtet werden, denn Threads, die lange aufeinander warten müssen oder sich häufig gegenseitig blockieren, verringern den Vorteil der Parallelisierung.

Für alle Aufgaben, die sich datenparallel abarbeiten lassen, so auch für die zuvor besprochenen Beispiele (Teilaufgaben: Bestimme die Flugrichtung von Raumschiff i; Bestimme die Laufrichtung von Fußgänger i; Bestimme die Marktpreise für Sonnensystem i etc.), bietet sich der Einsatz von OpenCL an. Als Programmierer müssen wir uns nicht darum kümmern, wie viele Prozessor-Threads oder Shader-Einheiten auf den jeweiligen Rechnern zur Verfügung stehen, denn für die Aufteilung der Berechnungen ist die OpenCL Runtime verantwortlich. Wir können uns also uneingeschränkt mit der Parallelisierung der eigentlichen Aufgabe befassen.

Asteroidenfeldsimulation

Zur Veranschaulichung der einzelnen Parallelisierungsmodelle finden Sie auf der Leser-CD unter anderem eine einfache OpenCL-basierte Asteroidenfeld-Bewegungssimulation (Abb. 2). Die einzelnen Asteroiden bewegen sich hierbei in zufällig festgelegte Richtungen innerhalb eines Simulationsbereichs mit begrenzter Ausdehnung, wobei es hin und wieder zu Kollisionen zwischen einzelnen Asteroiden kommt. Bevor wir nun die einzelnen Berechnungen parallelisieren können, müssen wir die Simulation zunächst in geeignete Teilschritte zerlegen:

Abb. 2: OpenCL-basierte Simulation eines Asteroidenfelds

  1. Positionen der einzelnen Asteroiden mitteln. Der Mittelpunkt markiert den Ursprung des verwendeten Kollisionsgitters.
  2. Bewegungssimulation der einzelnen Asteroiden.
  3. Berechnung der Kollisionsgitterpositionen und -indices (Collision Hash Grid) aller Asteroiden für eine effiziente Kollisionsprüfung.
  4. Durchführung der Kollisionsberechnungen (nur Asteroiden mit gleichem Gitterindex können miteinander kollidieren).

Da wir im Rahmen des hier betrachteten Programmbeispiels ganz bewusst auf die Simulation von Gravitationskräften verzichten, lassen sich die Bewegungen der einzelnen Asteroiden samt der zugehörigen Kollisionsgitterpositionen und -indices völlig unabhängig vom Verhalten der anderen Asteroiden in parallelen Threads berechnen. Für die Aufgaben 2 und 3 bietet sich infolgedessen der Einsatz von OpenCL an. Die für beide Aufgaben erforderlichen Berechnungen lassen sich hierbei nacheinander in einem einzigen OpenCL-Programm (Kernel) durchführen. Den grundlegenden Aufbau eines solchen Programms können Sie anhand von Listing 1 nachvollziehen.


__kernel void MasspointMovement(__global float4* Position,
                                __global float4* Velocity,
                                const float TimeStep)
{
    // ID des Asteroiden abfragen (global work-item ID):
    size_t id = get_global_id(0);
    Position[id] += Velocity[id]*TimeStep;
}

Im Zuge der Parallelisierung werden die einzelnen Berechnungen auf so genannte Work Items aufgeteilt. Im vorliegenden Fall entspricht jedes Work Item einem separaten Asteroiden. Um die neue Position eines Asteroiden berechnen zu können, müssen wir lediglich zuvor die Work-Item-ID beziehungsweise die Asteroiden-ID mithilfe der get_global_id-Funktion abfragen. Auf einer Vier-Kern-CPU mit acht Prozessor-Threads verteilt die OpenCL Runtime alle abzuarbeitenden Work Items möglichst gleichmäßig auf acht Worker Threads (Kernel-Instanzen). Ein vereinfachter Ablauf des datenparallelen Multithreadings ist in Listing 2 skizziert. Aufgeteilt werden die einzelnen Work Items gemäß Listing 2 nach folgendem Schema:

  • Thread 1: verantwortlich für die Work Items: 0, 8, 16, 24, …
  • Thread 2: verantwortlich für die Work Items: 1, 9, 17, …
  • […]
  • Thread 8: verantwortlich für die Work Items: 7, 15, 23, …

Im Unterschied zu dieser vereinfachten Darstellung werden in OpenCL darüber hinaus einzelne Work Items zu so genannten Work Groups zusammengefasst, wodurch bei Bedarf eine Synchronisierung der an einer Work Group beteiligten Work Items ermöglicht wird.

unsigned int WINAPI WorkerThread(void* data)
{
    [...]
    for(;;)
    {
        [...]
        for(WorkItemsID = ThreadID; WorkItemsID < NumWorkItems;
            WorkItemsID += NumThreads)
        {
            KernelFunction(WorkItemsID);
        }
        [...]
    }
    return 0;
}

Teilschritt 4, die Erkennung bevorstehender Kollisionen und die Berechnung der Geschwindigkeitsänderungen während dieser Kollisionen droht bei einer großen Anzahl von Objekten schnell zum Flaschenhals eines jeden Spiels zu werden – selbst bei Verwendung von Kollisionsgittern oder anderer Techniken (Beispiel: Sweep and Prune) zum Ausschluss einer großen Anzahl von Kollisionspaarungen. Zur Erklärung: Die Anzahl aller möglichen Kollisionen zwischen N Objekten berechnet sich wie folgt:

N Objekte: N*(N-1)/2 mögliche Kollisionspaare

Bsp.: 100 Objekte: 4950 mögliche Kollisionen

Streng genommen dürften die Kollisionsberechnungen nicht unabhängig von den Bewegungsberechnungen durchgeführt werden, da jede Kollision die Geschwindigkeiten und damit auch die Bewegungsrichtungen der an einer Kollision beteiligten Objekte verändert. Würde man die Berechnungen auf mehrere Threads aufteilen (z. B. ein Thread für die Bewegungssimulation sowie ein Thread für die Kollisionsberechnungen), so müssten die anstehenden Kollisionsberechnungen für ein Objekt so lange blockiert werden, bis die laufenden Bewegungsberechnungen abgeschlossen sind. In der Praxis (Programmbeispiele auf der Leser-CD) zeigt sich jedoch, dass sich Kollisionsberechnungen auch asynchron in einem separaten Thread durchführen lassen. Verzichtet man auf die Mittel der Synchronisierung, bewegen sich die kollidierenden Objekte im schlimmsten Fall für die Dauer der Kollisionsberechnungen mit der falschen Geschwindigkeit, was jedoch kaum einem Spieler jemals auffallen wird. Selbst wenn sich die Berechnungen bei einer großen Anzahl von Objekten über mehrere Frames hinziehen würden, stellt das in der Regel kein Problem dar, sofern sich die Objekte nicht allzu schnell bewegen.

[ header = Initialisierung einer OpenCL-Anwendung]

Initialisierung einer OpenCL-Anwendung

Die oberste Instanz jeder OpenCL-Anwendung wird als Host bezeichnet. Zu den Aufgaben einer Host-Anwendung zählt die Speicherverwaltung, die Auswahl der zu verwendenden OpenCL-Plattform (AMD, Intel, NVIDIA) und Geräte (Devices: CPU, GPU) sowie die Initialisierung von Gerätekontexten (Device Contexts) und Befehlswarteschlangen (Command Queues) zum Ansteuern der ausgewählten Geräte. Command Queues dienen zur Koordinierung alle zeitlichen Abläufe (Beispiele: Daten in einen Buffer schreiben bzw. aus einem Buffer auslesen, Ausführung eines OpenCL-Programms etc.). Standardmäßig werden die einzelnen Befehle nacheinander abgearbeitet, wobei eine neue Aktion erst dann gestartet wird, nachdem die vorangegangene beendet ist. Sollen die Befehle jedoch zum frühestmöglichen Zeitpunkt abgearbeitet werden, muss ein Command Queue mit dem CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE-Flag initialisiert werden. Die Schnittstelle zwischen der Host-Anwendung und einem OpenCL Device wird als Gerätekontext (Device Context) bezeichnet. Der Gerätekontext ist verantwortlich für die Ausführung der OpenCL-Programme, für die Verarbeitung der an einen Command Queue übermittelten Befehle und für die Verwaltung des verfügbaren Device-Speichers (z. B. GPU RAM). Die Auswahl einer passenden Plattform und der zu verwendenden Geräte sowie die sich daran anschließende Initialisierung eines Gerätekontexts und Command Queues lassen sich anhand von Listing 3 nachvollziehen.

cl_platform_id Platform = NULL;
cl_device_id Device = NULL;
cl_context DeviceContext = NULL;
cl_command_queue DeviceCommandQueue = NULL;

[...]

// erste gefundene Plattform auswählen:
cl_int error = clGetPlatformIDs(1, &Platform, NULL);

// GPU-Device verwenden:
error = clGetDeviceIDs(Platform, CL_DEVICE_TYPE_GPU, 1, &Device, NULL);

// Gerätekontext erzeugen:
DeviceContext = clCreateContext(NULL, 1, &Device, NULL, NULL, &error);

// Command Queue erzeugen:
DeviceCommandQueue = clCreateCommandQueue(DeviceContext, Device,
                     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);

Speicherverwaltung

Der von einer OpenCL-Anwendung verwendete Speicher ist gemäß Abbildung 1 in fünf verschiedene Bereiche unterteilt. Alle Daten, die zur Verarbeitung an ein Gerät (z. B. an die Grafikkarte) versandt werden sollen, befinden sich zunächst im Host-Speicher (Host Memory), dem regulären Arbeitsspeicher des Hauptprogramms. Der auf einem OpenCL-Gerät verfügbare Speicher ist in vier verschiedene Adressräume gegliedert. Gekennzeichnet werden die unterschiedlichen Adressräume durch die folgenden Address Space Qualifier:

  • __global (global)
  • __constant (constant)
  • __local (local)
  • __private (private)

Der konstante Speicher (Constant Memory) ist für alle Daten reserviert, die während der Berechnungen unverändert bleiben sollen. Im Unterschied zur Host-Anwendung besitzen die einzelnen Kernel-Instanzen lediglich Lesezugriff auf diesen Speicherbereich. Die übrigen Daten, auf die alle Instanzen sowohl Lese- als auch Schreibzugriff haben sollen, werden im globalen Speicher (Global Memory) abgelegt. Im Gegensatz zum globalen Speicher haben auf den lokalen Speicherbereich (Local Memory) lediglich die Kernel-Instanzen der gleichen Work-Group Zugriff. Schreib- und Lesezugriffe auf den lokalen Speicher sind um ein Vielfaches schneller als auf den globalen und konstanten Speicher. Variablen, die innerhalb einer Kernel-Funktion deklariert werden (oder innerhalb einer Funktion, die ihrerseits von einer Kernel-Funktion aufgerufen wird), befinden sich automatisch im privaten Speicherbereich.

Programm- und Kernel-Objekte

Verantwortlich für die durchzuführenden Berechnungen sind die in der Sprache OpenCL C geschriebenen Programme, die so genannten Kernels. Normalerweise werden OpenCL-Programme zur Laufzeit speziell für ein zuvor ausgewähltes Gerät kompiliert (Online Compilation). Bei Bedarf lassen sich jedoch auch bereits im Vorfeld kompilierte Programme verwenden (Offline Compilation). In der Regel ist dieses Vorgehen jedoch nicht zu empfehlen, da ein OpenCL-Programm hierbei für jedes denkbare Gerät separat kompiliert und gespeichert werden muss. Für die Verwaltung des Programmcodes samt der kompilierten Version ist das Programmobjekt cl_program zuständig. Das Kernel-Objekt cl_kernel ist hingegen für die Programmausführung auf dem zuvor gewählten OpenCL-Gerät verantwortlich. Die Initialisierung von Programm- und Kernel-Objekten wird in Listing 4 demonstriert.

cl_program Program = NULL;
cl_kernel  Kernel  = NULL;

static const char BuildOption__cl_fast_relaxed_math[] = "-cl-fast-relaxed-math";

[...]

// Kernel-Quellcode aus einer Textdatei einlesen:
char* Data = NULL;
size_t src_size;

Data = ReadCLFile(pFileName, &src_size);

const char * source = Data;

cl_int error;
Program = clCreateProgramWithSource(DeviceContext, 1,
                                    &source, &src_size, &error);
free(Data);

// Quellcode kompilieren:
if(fastMath == true)
    error = clBuildProgram(Program, 1, &Device,
            BuildOption__cl_fast_relaxed_math, NULL, NULL);
else
    error = clBuildProgram(Program, 1, &Device, NULL,
            NULL, NULL);

// Build-Informationen (Fehler, Warnungen, etc.) ausgeben:
char* build_log;
size_t log_size;

// Schritt 1 – Umfang der Build-Informationen ermitteln:
clGetProgramBuildInfo(Program, Device,
                      CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);

// Schritt 2 - Build-Informationen zwischenspeichern:
build_log = new char[log_size+1];

clGetProgramBuildInfo(Program, Device,
                      CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);

build_log[log_size] = '';

// Schritt 3: Build-Informationen in ein Logfile schreiben:
Add_To_Log(build_log);
delete[] build_log;

// Kernel-Objekt erzeugen:
Kernel = clCreateKernel(Program, pKernelName, &error);

Memory-Objekte

Memory-Objekte (cl_mem) ermöglichen den Zugriff auf Teile des globalen und konstanten Speichers eines OpenCL-Geräts. Entsprechend ihres Verwendungszwecks unterscheidet man zwischen Buffer- und Image-Objekten. Buffer-Objekte werden zum Speichern von eindimensionalen Datensätzen (z. B. eine Liste von float-Werten, eine Liste von vierdimensionalen float-Vektoren etc.) verwendet, Image-Objekte zum Speichern von zwei- oder dreidimensionalen Bildern. Die Initialisierung eines Buffer-Objekts können Sie anhand von Listing 5 nachvollziehen. Memory-Objekte, deren Daten in einem Kernel nicht verändert werden dürfen, sollten mit dem Flag CL_MEM_READ_ONLY initialisiert werden. Für Memory-Objekte, die lediglich zum Speichern der Kernel-Berechnungen dienen, ist das Flag CL_MEM_WRITE_ONLY vorgesehen. Das CL_MEM_READ_WRITE-Flag erlaubt sowohl Lese- als auch Schreibzugriffe.

cl_mem OpenCLBuffer = NULL;

[...]

// Buffer erzeugen (Beispiel: Array von vierdimensionalen float-Vektoren):
size_t BufferSize = 4*sizeof(float)*NumArrayElementsMax;

// Für Berechnungen auf der CPU immer das CL_MEM_USE_HOST_PTR-Flag verwenden!
if(OpenCLBuffer_UseHostPtr == true)
{
    OpenCLBuffer = clCreateBuffer(DeviceContext,
                   CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, BufferSize,
                   DataArray, NULL);
}
else
{
    OpenCLBuffer = clCreateBuffer(DeviceContext,
                   CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, BufferSize,
                   DataArray, NULL);
}

Wie bereits zuvor besprochen, müssen alle Daten, die von einem OpenCL-Gerät verarbeitet werden sollen, zunächst vom Host-Speicher in den globalen Gerätespeicher (in ein Buffer- oder Image-Objekt) kopiert werden. Sind die Berechnungen abgeschlossen, müssen die Ergebnisse anschließend in den Host-Speicher zurückkopiert werden. Bei Verwendung von Buffer-Objekten kommen hierbei die OpenCL-Funktionen clEnqueueWriteBuffer sowie clEnqueueReadBuffer zum Einsatz. Bei CPU-Berechnungen ist die Unterscheidung zwischen Host-Speicher und globalem Speicher genau genommen unsinnig, ebenso wie die zuvor beschriebenen Kopiervorgänge. Um einem OpenCL-Programm den direkten Zugriff auf den Host-Speicher zu ermöglichen (Memory Mapping, zu Deutsch Speichereinblendung), sollten Memory-Objekte, die lediglich für CPU-Berechnungen genutzt werden, mit dem CL_MEM_USE_HOST_PTR-Flag initialisert werden. Im Unterschied dazu sollte das CL_MEM_COPY_HOST_PTR-Flag nur für Memory-Objekte verwandt werden, die beispielsweise bei GPU-Berechnungen zum Einsatz kommen. Wie sich die Daten in einen Buffer schreiben beziehungsweise wieder auslesen lassen, wird in den Listings 6 und 7 gezeigt.

// Beispiel: Array von vierdimensionalen float-Vektoren:
size_t BufferSize = 4*sizeof(float)*NumElementsUsed;

if(OpenCLBuffer_UseHostPtr == false)
clEnqueueWriteBuffer(DeviceCommandQueue, OpenCLBuffer,
                     CL_TRUE, 0, BufferSize, DataArray, 0, NULL, NULL);
else
{
DataArray = (float*)clEnqueueMapBuffer(DeviceCommandQueue,
       OpenCLBuffer, CL_TRUE, CL_MAP_WRITE, 0, BufferSize, 0, NULL, NULL, NULL);

clEnqueueUnmapMemObject(DeviceCommandQueue, OpenCLBuffer,
                        DataArray, 0, NULL, NULL);
}
// Beispiel: Array von vierdimensionalen float-Vektoren:
size_t BufferSize = 4*sizeof(float)*NumElementsUsed;

if(OpenCLBuffer_UseHostPtr == false)
    clEnqueueReadBuffer(DeviceCommandQueue, OpenCLBuffer,
                        CL_TRUE, 0, BufferSize, DataArray, 0, NULL, NULL);
else
{
DataArray = (float*)clEnqueueMapBuffer(DeviceCommandQueue,
        OpenCLBuffer, CL_TRUE, CL_MAP_READ, 0, BufferSize, 0, NULL, NULL, NULL);

clEnqueueUnmapMemObject(DeviceCommandQueue,  OpenCLBuffer,
                        DataArray, 0, NULL, NULL);
}

Parameterübergabe an eine Kernel-Funktion

Nachdem die benötigten Kernel- und Memory-Objekte initialisiert worden sind, müssen im nächsten Schritt alle für die Berechnung erforderlichen Parameter mithilfe der clSetKernelArg-Funktion an die Kernel-Instanzen übergeben werden. Die in Listing 1 vorgestellte Kernel-Funktion erwartet beispielsweise drei Parameter: zwei Buffer-Objekte (Position und Velocity) sowie einen einfachen float-Parameter (TimeStep):

clSetKernelArg(Kernel, 0 /* erster Kernel Parameter*/,
               sizeof(Position), &Position);
clSetKernelArg(Kernel, 1 /* zweiter Kernel Parameter*/,
               sizeof(Velocity), &Velocity);
clSetKernelArg(Kernel, 2 /* dritter Kernel Parameter*/,
               sizeof(float), &TimeStep);

Durchführen von Berechnungen

Für die Durchführung der Berechnungen ist die clEnqueueNDRangeKernel-Funktion verantwortlich. Im einfachsten Fall, wenn auf die Synchronisierung von mehreren Berechnungen verzichtet werden kann und man es der OpenCL Runtime überlassen möchte, die einzelnen Work Items zu passenden Work Groups zusammenzufassen, gestaltet sich der Aufruf der clEnqueueNDRangeKernel-Funktion wie folgt:

clEnqueueNDRangeKernel(DeviceCommandQueue, Kernel,
                       WorkGroupDimension /* 1, 2 oder 3*/, NULL,
                       &MaxBufferElements /*Anzahl der Work-Items*/,
                       NULL, 0, NULL, NULL);

Über den Parameter WorkGroupDimension lässt sich festlegen, wie die OpenCL Runtime den Index eines Work Items in einer Work Group ermitteln soll. Verarbeitet eine Kernel-Funktion beispielsweise Daten in Form von Buffer-Objekten (Listing 1), so ist die Work-Group-Dimension 1 zu verwenden. Die Dimensionen 2 und 3 kommen hingegen bei der Verarbeitung von zwei- beziehungsweise dreidimensionalen Image-Objekten zum Einsatz.

Synchronisieren von Berechnungen

Die Synchronisation mehrerer Berechnungsschritte erfolgt mithilfe von Ereignissen (cl_event). In den Listings 8 und 9 sind zwei Möglichkeiten beschrieben, wie sichergestellt werden kann, dass die von den Kernel-Objekten Kernel_1 und Kernel_2 durchgeführten Berechnungsschritte auch wirklich nacheinander erfolgen. In Listing 8 wird der weitere Programmablauf im Host Thread mithilfe der clWaitForEvents-Funktion so lange unterbrochen, bis Kernel_1 seine Berechnungen abgeschlossen hat. Die explizite Verwendung besagter Funktion ist jedoch eigentlich nicht erforderlich, denn die zu komplettierenden Ereignisse lassen sich auch direkt als Parameter an die clEnqueueNDRangeKernel-Funktion übergeben (Listing 9).

cl_event event1;

clEnqueueNDRangeKernel(DeviceCommandQueue, Kernel_1,
WorkGroupDim, NULL, &MaxBufferElements, NULL, 0, NULL, event1);

clWaitForEvents(1, &event1);
clReleaseEvent(event1);

clEnqueueNDRangeKernel(DeviceCommandQueue, Kernel_2,
WorkGroupDim, NULL, &MaxBufferElements, NULL, 0, NULL, NULL);
cl_event event1;

clEnqueueNDRangeKernel(DeviceCommandQueue, Kernel_1,
WorkGroupDim, NULL, &MaxBufferElements, NULL, 0, NULL, event1);

clEnqueueNDRangeKernel(DeviceCommandQueue, Kernel_2,
WorkGroupDim, NULL, &MaxBufferElements, NULL, 1, event1, NULL);

clReleaseEvent(event1);
Unsere Redaktion empfiehlt:

Relevante Beiträge

Meinungen zu diesem Beitrag

X
- Gib Deinen Standort ein -
- or -