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

Kommentare

Infos zur Publikation

LU 02/2017: SICHER IM NETZ

Digitale Ausgabe: Preis € 5,95
(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

Samba-Server für Win-Daten & Linux-Dateirechte
Gerd Grundmann, 12.01.2017 18:25, 0 Antworten
Hallo Gemeinde, ich habe ewig gegooo.. aber nicht wirklich gefunden, was ich suche. Auf mei...
LM Cinnamon 18.1Serena 64 bit
Herwig Ameisbichler, 11.01.2017 13:02, 6 Antworten
Hallo an alle. Bin Linux Neuling.Habe auf meiner SSD neben W10 (finde ich nicht mehr)LM17.2 auch...
Bilder in E-Mail einbinden?
Werner Hahn, 23.12.2016 12:18, 4 Antworten
Dell Latitude E6510, Ubuntu 16.04, Mozilla Firefox Ich bitte um eine leicht verständliche Anle...
Linux Mint 17.2
Herwig Ameisbichler, 17.12.2016 11:19, 4 Antworten
Wer kann mir helfen? Hab seit gestern Linux am Pc.Wie kann ich diese Miniaturschrift im Menü,Lei...
Linux 17.2 Treiber f. PC Drucker installieren
Ernst Malitzki, 11.12.2016 13:04, 7 Antworten
Habe einen PC Drucker von HP. Möchte diesen PC Drucker an meinen PC anschließen. Hierzu benötige...