AA_dragster_123rf-9638361_SteveMann.jpg

© Steve Mann, 123rf.com

Gib Gas!

OpenCL-Workshop, Teil 2: Praxis

19.07.2011
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
 

Ähnliche Artikel

Kommentare

Infos zur Publikation

LU 11/2014: VIDEOS BEARBEITEN

Digitale Ausgabe: Preis € 4,95
(inkl. 19% MwSt.)

Mit der Zeitschrift LinuxUser sind Sie als Power-User, Shell-Guru oder Administrator im kleinen Unternehmen monatlich auf dem aktuelle Stand in Sachen Linux und Open Source.

Sie sind sich nicht sicher, ob die Themen Ihnen liegen? Im Probeabo erhalten Sie drei Ausgaben zum reduzierten Preis. Einzelhefte, Abonnements sowie digitale Ausgaben erwerben Sie ganz einfach in unserem Online-Shop.

NEU: DIGITALE AUSGABEN FÜR TABLET & SMARTPHONE

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

Tipp der Woche

Schnell Multi-Boot-Medien mit MultiCD erstellen
Schnell Multi-Boot-Medien mit MultiCD erstellen
Tim Schürmann, 24.06.2014 12:40, 0 Kommentare

Wer mehrere nützliche Live-Systeme auf eine DVD brennen möchte, kommt mit den Startmedienerstellern der Distributionen nicht besonders weit: Diese ...

Aktuelle Fragen

WLAN-Signalqualität vom Treiber abhängig
GoaSkin , 29.10.2014 14:16, 0 Antworten
Hallo, für einen WLAN-Stick mit Ralink 2870 Chipsatz gibt es einen Treiber von Ralink sowie (m...
Artikelsuche
Erwin Ruitenberg, 09.10.2014 07:51, 1 Antworten
Ich habe seit einige Jahre ein Dugisub LinuxUser. Dann weiß ich das irgendwann ein bestimmtes Art...
Windows 8 startet nur mit externer Festplatte
Anne La, 10.09.2014 17:25, 6 Antworten
Hallo Leute, also, ich bin auf folgendes Problem gestoßen: Ich habe Ubuntu 14.04 auf meiner...
Videoüberwachung mit Zoneminder
Heinz Becker, 10.08.2014 17:57, 0 Antworten
Hallo, ich habe den ZONEMINDER erfolgreich installiert. Das Bild erscheint jedoch nicht,...
internes Wlan und USB-Wlan-Srick
Gerhard Blobner, 04.08.2014 15:20, 2 Antworten
Hallo Linux-Forum: ich bin ein neuer Linux-User (ca. 25 Jahre Windows) und bin von WIN 8 auf Mint...