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 09/2015: Paketverwaltung

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

Grammatikprüfung in LibreOffice nachrüsten
Grammatikprüfung in LibreOffice nachrüsten
Tim Schürmann, 24.04.2015 19:36, 2 Kommentare

LibreOffice kommt zwar mit einer deutschen Rechtschreibprüfung und einem guten Thesaurus, eine Grammatikprüfung fehlt jedoch. In ältere 32-Bit-Versionen ...

Aktuelle Fragen

Scanner EPSON Perfection V 300 photo und VueScan
Roland Welcker, 19.08.2015 09:04, 1 Antworten
Verehrte Linux-Freunde, ich habe VueScan in /usr/local/src/vuex_3295/VueScan installiert, dazu d...
Empfehlung gesucht Welche Dist als Wirt für VM ?
Roland Fischer, 31.07.2015 20:53, 2 Antworten
Wer kann mir Empfehlungen geben welche Distribution gut geeignet ist als Wirt für eine VM für Win...
Plugins bei OPERA - Linux Mint 17.1
Christoph-J. Walter, 23.07.2015 08:32, 3 Antworten
Beim Versuch Video-Sequenzen an zu schauen kommt die Meldung -Plug-ins und Shockwave abgestürzt-....
Wird Windows 10 update/upgrade mein Grub zerstören ?
daniel s, 22.07.2015 08:31, 7 Antworten
oder rührt Windows den Bootloader nicht an? das ist auch alles was Google mir nicht beantw...
Z FUER Y UND ANDERE EINGABEFEHLER AUF DER TASTATUR
heide marie voigt, 10.07.2015 13:53, 2 Antworten
BISHER konnte ich fehlerfrei schreiben ... nun ist einiges drucheinander geraten ... ich war bei...