AA_dragster_123rf-9638361_SteveMann.jpg

© Steve Mann, 123rf.com

OpenCL-Workshop, Teil 2: Praxis

Gib Gas!

OpenCL macht die Grafikkarte vom alltäglichen, aber langweiligen Begleiter zum Turbo für anspruchsvolle Rechenoperationen. Wie im richtigen Leben spielt der Lader seine Qualitäten aber erst auf der Autobahn aus – im Stadtverkehr rührt sich wenig.

Nachdem der erste Teil dieses Artikels in der letzten Ausgabe einen Überblick über den Anwendungsbereich von OpenCL gab und auf die Installation einging, widmen sich dieser Teil den Arbeitsschritten zur Bildfaltung mittels OpenCL auf der Grafikkarte (Abbildung 1, Abbildung 2). Die im vorigen Teil des Artikels vorgestellte CPU-Implementierung dient dabei später als Referenz für die Laufzeitmessung.

Abbildung 1: Das Eingabebild: Larry Ewings Tux.
Abbildung 2: Das Ergebnis: Tux' Gradienten, gefaltet mit einem Sobelkernel.

Nun widmen wir uns dem OpenCL-Code, dessen wichtigste Partien Sie in Listing 1 sehen. Den kompletten Quelltext finden Sie auch auf der Heft-DVD sowie zum Download auf unserer Website [1].

Listing 1

 1 #define __CL_ENABLE_EXCEPTIONS... 13 const char* kernelSource = "\ 14 __kernel void convolveKernel(\ 15          global uchar *in,\... 18          global uint *out,\... 21          global float *convKernel,\ 22          uint convKernelWidth,\ 23          uint convKernelHeight)\ 24 {\ 25  size_t x = get_global_id(0);\ 26  size_t y = get_global_id(1);\... 35   for(size_t ky = 0; ky < convKernelHeight; ++ky)\ 36    { 37     for(size_t kx = 0; kx < convKernelWidth; ++kx)\ 38      {\ 39       convolutionSum += (float) in[(y + ky) * inWidth + (x + kx)]\ 40                          * convKernel[ky * convKernelWidth + kx];\ 41      }\ 42   }\ 43  out[y * outWidth + x] = (uint) clamp(convolutionSum, 0, 255);\ 44 }"; 45 46 /** 47  * Convolve a grayscale image with a convolution kernel on the GPU using OpenCL. 48  */ 49 grayImage convolveGPU(grayImage in, convolutionKernel convKernel) 50 { 51  grayImage out; 52  out.width = in.width - (convKernel.width - 1); 53  out.height = in.height - (convKernel.height - 1); 54  out.data = new uchar[out.height * out.width]; 55 56  // Platforms 57  std::vector< cl::Platform > platforms; 58  cl::Platform::get(&platforms); 59  assert(platforms.size() > 0); 60 61  // Devices 62  std::vector<cl::Device> devices; 63  platforms[0].getDevices(CL_DEVICE_TYPE_GPU, &devices); 64  assert(devices.size() > 0); 65  assert(devices[0].getInfo<CL_DEVICE_TYPE>() == CL_DEVICE_TYPE_GPU); 66 67  // Context 68  cl::Context context(devices); 69 70  // Create GPU buffers 71  cl::Buffer inGPU(context, CL_MEM_READ_ONLY, in.width * in.height * sizeof(uchar)); 72  cl::Buffer convKernelGPU(context, CL_MEM_READ_ONLY, convKernel.width * convKernel.height * sizeof(float)); 73  cl::Buffer outGPU(context, CL_MEM_WRITE_ONLY, out.width * out.height * sizeof(uint)); 74 75  // Commandqueue 76  cl::CommandQueue queue(context, devices[0], 0); 77 78  // Upload in.data to inGPU 79  queue.enqueueWriteBuffer(inGPU, false, 0, in.width * in.height * sizeof(uchar), in.data);... 86  // Upload kernel.data to convKernelGPU 87  queue.enqueueWriteBuffer(convKernelGPU, true, 0, convKernel.width * convKernel.height * sizeof(float), convKernel.data);... 94  // Program 95  cl::Program::Sources source(1, std::make_pair(kernelSource, strlen(kernelSource))); 96 97  cl::Program program(context, source); 98  program.build(devices); 99100  // Ranges101  size_t localWidth = 16;102  size_t localHeight = 16;103104  cl::NDRange localRange(localWidth, localHeight);105  cl::NDRange globalRange(((out.width-1)/localWidth+1) * localWidth, ((out.height-1)/localHeight+1) * localHeight);106107  // Run kernel108  cl::Kernel kernel(program, "convolveKernel");109  cl::KernelFunctor func = kernel.bind(queue, globalRange, localRange);110111  cl::Event event = func(inGPU, in.width, in.height, outGPU, out.width, out.height, convKernelGPU, convKernel.width, convKernel.height);112  event.wait();113114  // Download result115  uint *outTemp = new uint[out.width * out.height];116  queue.enqueueReadBuffer(outGPU, true, 0, out.width * out.height * sizeof(uint), outTemp);...123  // Convert uint array to uchar array124  for(size_t i = 0; i < out.width * out.height; ++i)125   { out.data[i] = (uchar) outTemp[i];127  }...129  delete outTemp;130  return out;131 }

Los geht's

Um die in Teil 1 bereits erwähnten OpenCL-C++-Bindings zu verwenden, genügt ein simples #include <CL/cl.hpp>. Statt der in C üblichen Errorcodes sollen die Exceptions der Bindings zum Einsatz kommen. Dazu dient die Definition #define __CL_ENABLE_EXCEPTIONS. Sämtliche Klassen finden sich im Namespace cl::. Um dem Linker mitzuteilen, wogegen er zu linken hat, fügen Sie der Parameterliste beim Aufruf von g++ noch das Argument -lOpenCL hinzu.

Benötigte Verwaltungsobjekte

Im Gegensatz zu Cuda erzeugt OpenCL den plattformabhängigen Code erst zur Laufzeit. Somit muss OpenCL zunächst herausfinden, welche Hardware den parallel auszuführenden Code abarbeiten wird. Dazu erzeugen Sie in den Zeilen 57 und 62 von Listing 1 eine cl::Platform und einen Vektor mit cl::Devices.

TIPP

Es gibt zwei unterschiedliche Typen von cl::Platform: "full profile" und "embedded profile". Dieser Artikel behandelt das "full profile".

Innerhalb jeder cl::Platform kann es mehrere cl::Device geben, die jeweils eine verbaute GPU oder CPU wiederspiegeln. Die assert()>-Anweisungen ab Zeile 64 stellen sicher, dass es mindestens ein OpenCL-fähiges Device gibt und es sich beim ersten Device um eine unterstützte Grafikkarte handelt. Für ein Ausführen auf der CPU müssen Sie dementsprechend dasjenige Device aus dem Vektor nehmen, das vom Typ CL_DEVICE_TYPE_CPU ist.

Ein cl::Context verwaltet über diverse cl::Devices hinweg Objekte wie Command Queues, Speicherobjekte, Kernel und Ausführungsobjekte. Auf diese Objekte werden wir im Folgenden noch eingehen; im Beispiel verwaltet ein cl::Context nur unsere Grafikkarte.

Schließlich gibt es noch eine cl::CommandQueue. In diese reihen Sie Aktionsobjekte ein, die das System dann im Standardfall der Reihe nach (FIFO) ausführt.

LinuxCommunity kaufen

Einzelne Ausgabe
 
Abonnements
 
TABLET & SMARTPHONE APPS
Bald erhältlich
Get it on Google Play

Deutschland

Ähnliche Artikel

  • OpenCL-Workshop, Teil 1: Grundlagen
    Grafikkarten können nicht nur bunte Bildchen malen: Beim parallelen Verarbeiten großer Datenmengen laufen die GPUs den CPUs den Rang ab. Dabei dient OpenCL unabhängig von Hardware und Hersteller als Programmierplattform.
Kommentare

Infos zur Publikation

LU 12/2016: Neue Desktops

Digitale Ausgabe: Preis € 5,99
(inkl. 19% MwSt.)

LinuxUser erscheint monatlich und kostet 5,95 Euro (mit DVD 8,50 Euro). Weitere Infos zum Heft finden Sie auf der Homepage.

Das Jahresabo kostet ab 86,70 Euro. Details dazu finden Sie im Computec-Shop. Im Probeabo erhalten Sie zudem drei Ausgaben zum reduzierten Preis.

Bei Google Play finden Sie digitale Ausgaben für Tablet & Smartphone.

HINWEIS ZU PAYPAL: Die Zahlung ist ohne eigenes Paypal-Konto ganz einfach per Kreditkarte oder Lastschrift möglich!

Aktuelle Fragen

Drucker Epson XP-332 unter ubuntu 14.04 einrichten
Andrea Wagenblast, 30.11.2016 22:07, 1 Antworten
Hallo, habe vergeblich versucht mein Multifunktionsgerät Epson XP-332 als neuen Drucker unter...
Apricity Gnome unter Win 10 via VirtualBox
André Driesel, 30.11.2016 06:28, 2 Antworten
Halo Leute, ich versuche hier schon seit mehreren Tagen Apricity OS Gnome via VirtualBox zum l...
EYE of Gnome
FRank Schubert, 15.11.2016 20:06, 2 Antworten
Hallo, EOG öffnet Fotos nur in der Größenordnung 4000 × 3000 Pixel. Größere Fotos werden nic...
Kamera mit Notebook koppeln
Karl Spiegel, 12.11.2016 15:02, 2 Antworten
Hi, Fotografen ich werde eine SONY alpha 77ii bekommen, und möchte die LifeView-Möglichkeit nu...
Linux auf externe SSD installieren
Roland Seidl, 28.10.2016 20:44, 1 Antworten
Bin mit einem Mac unterwegs. Mac Mini 2012 i7. Würde gerne Linux parallel betreiben. Aber auf e...