OpenCL - Open Computing Language

Dipl. Math. F. Braun Universität Regensburg  Rechenzentrum http://www.uniregensburg.de/EDV/kurs_info/brf09510/hpc/opencl/opencl.html http://www.uniregensburg.de/EDV/kurs_info/brf09510/hpc/opencl/opencl.pdf 19. Dezember 2016

KAPITEL 1

Einführung in OpenCL

1. Kursinhalt und Kursbeispiele 2. Designer Apple, danach Khronos; in Zusammenarbeit mit AMD, IBM, Intel, NVidia. Zitat von der Homepage von Khronos:

The Khronos Group was founded in January 2000 by a number of leading media-centric companies, including 3Dlabs, ATI, Discreet, Evans & Sutherland, Intel, NVIDIA, SGI and Sun Microsystems, dedicated to creating open standard APIs to enable the authoring and playback of rich media on a wide variety of platforms and devices. The Khronos Group is a not for prot, member-funded consortium focused on the creation of royalty-free open standards for parallel computing, graphics and dynamic media on a wide variety of platforms and devices. All Khronos members are able to contribute to the development of Khronos API specications, are empowered to vote at various stages before public deployment, and are able to accelerate the delivery of their cutting-edge 3D platforms and applications through early access to specication drafts and conformance tests.1

3. Literatur Literatur zum High performance computing:

http://www.uni-regensburg.de/EDV/kurs_info/brf09510/hpc/Literatur.html Aaftab Munshi, Benedict Gaster, Timothy G. Mattson, James Fung, Dan Ginsburg: OpenCL Programming Guide, 2012, 978-0-321-74964-2 Beispiele zu diesem Buch:

https://github.com/bgaster/opencl-book-samples OpenCL Seite:

https://www.khronos.org/opencl/ OpenCL Spezikationen:

https://www.khronos.org/registry/cl/ OpenCL Spezikation 2.1:

https://www.khronos.org/registry/cl/specs/opencl-2.1.pdf OpenCL Reference Card:

https://www.khronos.org/files/opencl21-reference-guide.pdf OpenCL bei Nvidia:

https://developer.nvidia.com/opencl OpenCL bei AMD:

http://developer.amd.com/tools-and-sdks/opencl-zone/ 1https://www.khronos.org/about,

22.1.2016

3

4

1. EINFÜHRUNG IN OPENCL

OpenCL bei Intel:

https://software.intel.com/en-us/intel-opencl OpenCL bei IBM:

https://www.ibm.com/developerworks/community/groups/service/html/ communityview?communityUuid=80367538-d04a-47cb-9463-428643140bf1

4. Geschichte 12'2008 Beginn der Entwicklung von OpenCL 6.10.2009 OpenCL 1.0 1.6.2011 OpenCL 1.1: 3-component vectors Additional image formats Multiple hosts and devices Buer regi-

on operations Enhanced event-driven execution Additional OpenCL C built-ins Improved OpenGL data/event interop 14.11.2012 OpenCL 1.2: Device partitioning Separate compilation and linking Enhanced image support Builtin kernels / custom devices Enhanced DX and OpenGL Interop 17.10.2014 OpenCL 2.0: Shared Virtual Memory On-device dispatch Generic Address Space Enhanced Image Support C11 Atomics Pipes Android ICD 11'2015 OpenCL 2.1: SPIR-V in Core Subgroups into core Subgroup query operations clCloneKernel Lowlatency device timer queries 5'2016 OpenCL 2.2 provisional

Quelle: https://www.khronos.org/assets/uploads/developers/library/overview/opencl_overview.pdf

5. Ziele und Eigenschaften OpenCL ist für heterogene Systeme (ein oder mehrere CPUs, GPUs, andere Beschleuniger) konzipiert. OpenCL ist standardisiert. Grundlage der Sprache für heterogene Systeme ist C99 OpenCL ist für Datenparallelität (SIMD, SIMT). Architekturen dieses Typs sind GPUs, CPUs, Cell-Prozessoren. . . OpenCL ist für die Formulierung von Algorithmen, die ezient auf diesen Architekturen laufen. OpenCL ist portabel und lauähig auf verschiedenen Platformen. OpenCL deckt heterogene Architekturen ab OpenCL und C99 haben eine gemeinsame Untermenge; OpenCL ist aus C99 abgeleitet, abgemagert und erweitert. Die Implementierung von OpenCL erfolgt durch die Hersteller der Hardware. OpenCL ist wie MPI eine Library, die von normalem C oder C++ aufgerufen wird. Der OpenCL-Code wird vom Compiler nicht beachtet. OpenCL konkurriert einerseits direkt mit Cuda. In beiden Sprachen werden GPUs für aufwendigere Rechnungen mit sog. kernels benutzt. Ein weiterer lange etablierter Konkurrent liegt mit OpenMP für Multi-Thread/MultiCore-Programmierung vor. Hier wird die CPU mit lightweight threads zum Rechnen verwendet. Der Begri kernel hat ursprünglich den wichtigen algorithmischen Kern eines Algorithmus bezeichnet und steht heute für die nicht auf der Host-CPU laufenden Programmteile. Für die jeweiligen Unterschiede müssen viele Fakten in Betracht gezogen werden: Für Cuda gibt es nur einen Hersteller (nvidia) und nur ein device (nvidia Graphikkarten). Für OpenCL gibt es mehrere Hersteller (AMD, Nvidia, Apple, Intel, IBM, Altera, Portable OpenCL) und mehrere devices (GPUs - nvidia, Radeon; CPUs; Xeon Phi). Damit erönen sich weit mehr Einsatzfelder. Trotzdem kann ein konkretes Programm überall, auf einer eingeschränkten Menge von devices oder auf nur einem einzigen device laufen. Falls mehrere device-Typen möglich sind, kann die performance auf nur einem dieser Typen gut sein. Hochwertiger Code sollte auf guten Implementierungen portabel sein. Code-Ezienz muss nicht portabel sein. Derzeit bietet Cuda gegenüber OpenCL wesentlich mehr hochwertige Bibliotheken wie cuFFT, cuLA u.v.a.m.

7. OPENCL AN DER UNIVERSITÄT REGENSBURG

5

Die Ausführung von Code auf anderen Maschinen setzt bei Cuda nur eine NVidia-Graphikkarte voraus. Bei OpenCL muss die richtige Runtime-Environment auf dem Zielrechner installiert sein. Das macht das Bereitstellen von Code zum Download (als source oder auch als binary) im Internet wesentlich schwerer. Cuda kann vom Compiler übersetzt werden; der Binärcode ist dann auf Systemen mit einer Graphikkarte lauähig. OpenCL-kernels können nicht übersetzt werden, da erst auf der Zielmaschine die OpenCL-Platform bekannt ist. Ser kernel-Code wird als erst zur Laufzeit compiliert.

6. Compiler und Systeme Zur Übersetzung von Programmen mit OpenCL braucht man ein OpenCL-SDK. Es wird von den Herstellern der Hardware, die vom Programm benutzt werden soll, erstellt und geliefert. Ein OpenCL-Programm enthält normalen Binärcode. Es kann allerdings nur gestartet werden, auf der Startplatform wenigstens eine OpenCL-Laufzeitumgebung existiert. Mehrere Laufzeitumgebungen können koexistieren; sie tragen herstellerspezische Informationen in /etc/OpenCL/vendors/*.icd ein. Hier bezieht das Programm auch die Informationen über nutzbare kernel-Hardware. Für Nvidia-Graphikkarten ist OpenCL in der Cuda-Installation dabei. Eine CPU-Implementierung für OpenCL wird von AMD für AMD- und Intel-CPUs, sowie Radeon-Graphikkarten bereitgestellt:

http://developer.amd.com/tools-and-sdks/opencl-zone/amd-accelerated-parallel-processing-app-sdk/ Nach Auswahl der gewünschten Version und Akzeptanz des AMD License Agreements erhält man eine .tar.bz2Datei. Nach dem Auspacken startet man als Superuser (sudo!) die enthaltene sh-Datei. Nach einem reboot sollte AMD-OpenCL laufen. Mit AMD-Graphikkarten benötigt man vor der Installation erst den AMD-GPU-Treiber. Anleitung: http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2012/10/AMD_APP_SDK_InstallationNotes.pdf Nur für Intel-CPUs (Xeon Phi, Xeon, Core) bietet Intel eine Implementierung an.

https://software.intel.com/en-us/intel-opencl

Apple unterstützt OpenCL für Nvidia- und Radeongraphikkarten und CPUs auf MacOS. Ohne Anspruch auf Vollständigkeit: Altera, AMD, Apple, ARM, Catalyst, Creative Technology, IBM, Imagination, Intel, Mediatek, Mesa, Narvell, Nvidia, Pocl, Qualcomm, Samsung, Sony Mobile, ST, TI Texas Instruments, Vivante, Xilinx, ZiiLABS. Open Source: Beignet, FreeOCL, Mali, OCLGrind, POCL.

(http://www.iwocl.org/resources/opencl-implementations/ https://github.com/fakenmc/cf4ocl/wiki/OpenCL

7. OpenCL an der Universität Regensburg NVidia auf RZ411, RZ412, CH200  CH209, PC58685, Athene (queue=nvidia) AMD auf PC55556 als guest

KAPITEL 2

crash course

1. OpenCL und C99/C++ OpenCL beruht auf C99. Wegen der speziellen Bedürfnisse heterogenen Rechnens gibt es Erweiterungen von C99, aber auch Einschränkungen, manchmal nur um den Implementierern die Arbeit zu erleichtern. Einschränkungen gegenüber C99:

• • • • • • • • • • • • • • • • •

• • • • •

keine C99-Standard Header; Ersatz ist der Vorrat von built-in-Funktionen. Kernel-Zeigerargumente müssen mit __global, __constant, oder __local qualiziert werden. Nur __constant-Zeiger dürfen auf __constant-Zeiger zugewiesen werden. keine Zeiger auf Funktionen. keine Bitfelder. keine VLAs (variable length arrays). keine variadic templates (C++). weder Makros noch Funktionen mit variabler Argumentzahl (auÿer printf und enqueue_kernel). keine auto und register Speicherklassen. keine rekursiven Funktionen. Kernels sind immer Prozeduren mit void-Ergebnis. keine Kernel-Argumente mit den Typen bool, half, size_t, ptrdiff_t, intptr_t, uintptr_t, sowie Strukturen und Unions mit einem dieser Komponenten. keine Arithmetik für half; half ist nur Speicherformat. irreduzible Anweisungsgruppen (Bsp.: Sprünge in Schleifen) sind implementation-dened. const, restrict und volatile sind erlaubt, aber für images verboten. event_t darf kein kernel-Argument sein. The event type (event_t) cannot be used as the type of a kernel function argument. The event type cannot be used to declare a program scope variable. The event type cannot be used to declare a structure or union eld. The event type cannot be used with the __local, __constant and __global address space qualiers. The clk_event_t, ndrange_t and reserve_id_t types cannot be used as arguments to kernel functions that get enqueued from the host. The clk_event_t and reserve_id_t types cannot be declared in program scope. The behavior of applying the sizeof operator to the queue_t, clk_event_t, ndrange_t and reserve_id_t types is implementation-dened. Kernels enqueued by the host must continue to have their arguments that are a pointer to a type declared to point to a named address space. A function in an OpenCL program cannot be called main. An image type ( image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buer_t, or image1d_array_ can only be used as the type of a function argument. An image function argument cannot be modied. Elements of an image can only be accessed using the Built-in Image Functions. An image type cannot be used to declare a variable, a structure or union eld, an array of images, a pointer to an image, or the return type of a function. An image type cannot be used with the __global, __private, __local and __constant address space qualiers. The sampler type (sampler_t) can only be used as the type of a function argument or a variable declared in the program scope or the outermost scope of a kernel function. The behavior of a sampler variable declared in a non-outermost scope of a kernel function is implementation-dened. A sampler argument or variable cannot be modied. The sampler type cannot be used to declare a structure or union eld, an array of samplers, a pointer to a sampler, or the return type of a function. The sampler type cannot be used with the __local and __global address space qualiers.

Vorsicht ist geboten bei: 7

8

2. CRASH COURSE

• Die Speicherdarstellung, speziell die size in Bytes vieler Datentypen kann (und wird) zwischen host und device variieren. Zusätzlich ist sie implementation-dened. • endianess kann zwischen host und device variieren. OpenCL hat zusätzlich zu C99 Vektortypen, Workitems, Workgroups, Synchronisierung, neue qualier (global, local, constant, private), Images, neue built-ins (min, max, radians, sign). Kernels haben weitere Beschränkungen. Kernel-Argumente dürfen nur globale, konstante oder lokale Zeiger sein. Zeigerverkettungen (Zeiger auf Zeiger) sind nicht erlaubt. Verboten sind bool, half, size_t, ptrdiff_t, intptr_t, uintptr_t, even_t. Das Kernel-Ergebnis ist immer void. Kernels werden auf devices ausgeführt. Sie bestehen aus Quellcode in Dateien oder Strings und werden compiliert. Es existiert eine portable Zwischensprache für Vorkompilierungen (SPIR, Standard Portable Intermediate Representation) Device-management, Speichertransporte, Synchronisierungen sind Aufgabe des Host. Die C++-Schnittstelle ist lediglich ein Wrapper. Sie benutzt die C-Library.

2. Compilation OpenCL ist scheinbar nur eine Library. Programme werden folglich mit einem C/C++-Compiler übersetzt und gegen die OpenCL-Library gelinkt. In manchen Installationen muss man mit include- und lib-Pfaden nachhelfen.

gcc -std=c11 x.c -lOpenCL g++ -std=c++14 x.cc -lOpenCL Scheinbar. . . ? Das merkt man spätestens, ween der übersetzte und montierte Code beim Starten nicht oder nicht richtig läuft. OpenCL-Programme enthalten zusätzlich OpenCL-Code. Dieser OpenCL-Code muss natürlich ebenfalls compiliert werden. Das erfolgt während des Programmlaufs mit einem in die Library integrierten Compiler. Das OpenCL-Programm wird mit clCreateProgramWithSource, clCreateProgramWithBinary oder clCreateProgramWithBuiltinKernels erzeugt. Die Übersetzung erfolgt mit dem Aufruf clBuildProgram online oder mit clCompileProgram und clLinkProgram oine. Ein garbage-collector gibt nicht mehr benötigten Speicher wieder frei, wenn sein Referenzzähler 0 wird. Create erhöht der Referenzzähler eines Programms implizit, clRetainProgram macht dasselbe explizit. Mit clReleaseProgram wird der Referenzzähler vermindert und der Speicher seiner Freigabe nähergebracht. Nach der Compilierung muss noch ein kernel-Objekt erzeugt werden. Das erfolgt mit clCreateKernel. Für kernels gilt derselbe retain/release-Mechanismus. Zum Start wird der kernel mit Argumenten versorgt (clSetKernelArg) und in eine Warteschlange eingereiht (clEnqueueNDRangeKernel. ,,

3. C/C++-Header Die Library wird mit Header-Files eingebunden. Apple genehmigt sich hier eine Extrawurst:

#include #include #include (Apple)

4. Fehlerbehandlung Alle Funktionen liefern einen Fehlercode des Typs cl_int. Meist erfolgt dies als Funktionsergebnis, manchmal als Variablenreferenz im letzten Argument.

cl_int err; err = clX() result = clX(..., &err)

6. PRINZIPIELLER ABLAUF EINES OPENCL-PROGRAMMS

9

if (err) ... Benutzen Sie zur Bequemlichkeit die beiden Dateien openclx.h und openclx.c aus svn/hpc/trunk/opencl. Wegen der nichtidentischen Schnittstelle werden Funktionen mit error-result als

ocl_verify(clBuildProgram(prog, 0, NULL, "", NULL, NULL)); und Funktionen mit error-Parameter als

cl_kernel k_rot13=clCreateKernel(prog, "rot13", &error); ocl_verify_err(error); aufgerufen. In C++ können Fehler durch Exceptions abgefangen werden. Das muss vor der Header-Datei angegeben werden.

#define __CL_ENABLE_EXCEPTIONS #include #include

5. Abfragen In OpenCL-C sind sehr viele Abfragen zweistug. Nur wenn die benötigte Gröÿe des Speichers im voraus bekannt ist, kann man auf die Zweistugkeit verzichten.

err = clGetX (..., 0, NULL, &size); data = malloc (size); err = clGetX (..., size, data, NULL); In C++ erfolgt die Speicherverwaltung automatisch im Wrapper.

cl::X::get()

Platform List: clGetPlatformIDs (ocls21 p62) Platform Info: clGetPlatformInfo (ocls21 p62) Device List: clGetDeviceIDs (ocls21 p64) Device Info: clGetDeviceInfo (ocls21 p66)

6. Prinzipieller Ablauf eines OpenCL-Programms 1. Steuerung erfolgt durch die/eine CPU (den Host) 2. Analyse der verfügbaren Komponenten und ihrer Leistungen (der Platforms) 3. Entscheidung über die geeigneten kernels; Erzeugung dieser kernel-Objekte 4. Erzeugung der zugehörigen memory-Objekte 5. Start der kernels mit optimaler Parallelität; dazu werden command-queues verwendet; Synchronisierungen steuern den Ablauf 6. Ergebnisse

10

2. CRASH COURSE

7. Objekte Objekte in C sind globale Variable und Funktionen. Objekte in C++ sind struct- und class-Variable. Objekte in OpenCL sind alle notwendigen Informationseinheiten zur Kommunikation zwischen Host und Device. Es gibt Buer objekte, Image objekte, Pipe objekte, Memory objekte, Sampler objekte, Program objekte, Kernel objekte, Event objekte. Objekte müssen erzeugt werden. Bei der Erzeugung ist das Funktionsergebnis das neue Objekt. Der Fehlercode wird im letzten Zeigerargument gespeichert. In C++ wird ein Konstruktor verwendet.

o = clCreateX (..., &err) o = cl::X (..., &err) In C müssen Objekte explizit mit Release freigegeben werden. In C++ erfolgt die Freigabe automatisch im Destruktor.

clReleaseX (o)

8. Wichtige Objektbeispiele

Context: cl_context ct = clCreateContext (props, numDev, dev, NULL, NULL, &err); Command queue: cl_command_queue_ cq = clCreateCommandQueue (ct, dev [0], CL_QUEUE_OUT_OF_ORDER_EXEC_ Program: cl_program pg = clCreateProgramWithSource (ct, 1, (const char **) &srcString, NULL, &err); Kernel: cl_kernel kl = clCreateKernel (pg, "name", &err); Memory object: cl_mem bf = clCreateBuffer (ct, CL_MEM_WRITE_ONLY, sz*sizeof(float), NULL, &err);

9. OpenCL-Konzept Ein OpenCL-Programm läuft auf einem steuernden Hostrechner und einem oder mehreren zur Beschleunigung verwendbaren devices. Mindestens eine OpenCL-Implementierung muss verfügbar sein. Die verfügbaren Implementierungen heiÿen platforms. Die erste Frage eines OpenCL-Programms muss also die nach den verfügbaren platforms und deren Eigenschaften sein. clGetPlatformIDs Ähnlich muss man anschlieÿend nach den platform-spezischen devices und deren Eigenschaften fragen. Pro Platform sind ein oder mehrere devices mit möglicherweise verschiedenen Typen (CPU/GPU) möglich. clGetDeviceIDs Kennt man die gewünschte Platform und das gewünschte device, dann kann man einen (oder mehrere) Kontext erzeugen, in dessen Rahmen Rechnungen von sog. kernels ausgeführt werden können. Ein solcher Kontext wird erzeugt und nach Verwendung wieder gelöscht. Im Kontext werden ein oder mehrere devices bereitgestellt, Speicher alloziert, gefüllt oder ausgelesen und kernels ausgeführt. Alle Tätigkeiten werden in Warteschlangen zur Ausführung eingereiht. Im Kontect werden schlieÿlich auch Ereignisse verwaltet, mit denen der Abschluÿ con Kommandos getestet oder abgewartet werden kann. clCreateContext(FromType) clRetainContext clReleaseContext clGetContextInfo Ohne Warteschlange kommt man im Kontext nicht aus. Auch Warteschlangen werden erzeugt und nach Abschluÿ aller Arbeiter wieder gelöscht. Sie dienen zur synchronen oder asynchronen Abarbeitung verschiedener Kommandos (Speicherzugri, Kernelausführung. . . ) clCreateCommandQueueWithProperties clSetDefaultDeviceCommandQueue

13. DIE SPRACHE OPENCL-C

11

clRetainCommandQueue clReleaseCommandQueue clGetCommandQueueInfo Speichern von Daten erfolgt in Buerobjekten, die erzeugt, gefüllt, gelesen und gelöscht werden müssen. clCreateBuer clCreateSubBuer Zum Schreiben und Lesen werden Befehle in eine Warteschlange eingetragen clEnqueueReadBuer clEnqueueWriteBuer clEnqueueReadBuerRect clEnqueueWriteBuerRect clEnqueueCopyBuer clEnqueueCopyBuerRect clEnqueueFillBuer clEnqueueMap Buer clEnqueueBuer clEnqueueBuer clEnqueueBuer clBuer clBuer clBuer

10. Warteschlangen Kommandos müssen in eine Warteschlange eingereiht werden.

11. Events 12. Beispiel 13. Die Sprache OpenCL-C OpenCl-Programme sind C99-ähnliche Texte mit Einschränkungen und Erweiterungen. Sie werden entweder in Strings gespeichert; der C-Compiler analysiert diese Texte nicht. Sie können auch in .cl-Dateien gespeichert werden und vom C-Compiler in eine String-Variable gelesen werden. Auch die Verarbeitung vorübersetzter Programm ist möglich; das ist jedoch nicht portabel, die Vorübersetzungen können nicht portiert werden.

13.1. Kernels.

Kernels sind die Ausführungseinheiten für ein device.

13.2. Datentypen. Die Datentypen von OpenCL sind gegenüber C99 anders und erweitert deniert. Ihre Denition ist jedoch für alle Platformen verbindlich. Diese Denition spiegelt sich in den Datentypen des Hostprogramms (cl_int) wider. Nicht deniert ist die Endianess in OpenCL. Wenn zwischen Host und Device verschiedene Enidianess vorliegt, muss der Programmierer die Bytereihenfolge explizit ändern. Noch unbefriedigender ist die Enianess bei Kernelargumenten geregelt:

When the host and device have dierent endianness, the developer must ensure that kernel argument values are processed correctly. The implementation may or may not automatically convert endianness of kernel arguments. Developers should consult vendor documentation for guidance on how to handle kernel arguments in these

12

2. CRASH COURSE

situations.

(ocl21, p. 281) Im Klartext heiÿt das, dass Kernargumente implementierungsabhängig mal konvertiert werden können, von anderen Implementierungen aber auch nicht! Genau hier bräuchte man aber für heterogene Systeme klare Denitionen. Und als Komfort für den Programmierer eigentlich automatische Endianess-Konversionen! Tabelle 1.

Datentypen von OpenCL-C

Typ Opt. Info. Di zu C Host true false 1 0 bool _Bool char 8 bit cl_char uchar 8 bit cl_uchar short 16 bit cl_short ushort 16 bit cl_ushort int 32 bit eigener Type cl_int 32 bit uint cl_uint long 64 bit cl_long ulong 64 bit cl_ulong half IEEE-754-08 cl_half oat IEEE-754 cl_oat double opt. IEEE-754 cl_double sizeof result size_t ptrdi_t (void*) - (void*) intptr_t void* uintptr_t void* void

Vektor charn ucharn shortn ushortn intn uintn longn ulongn halfn oatn doublen

Die in C übliche Schreibweise mit unsigned (unsigned int statt uint) ist erlaubt. Die ganzzahligen Typen sind eher wie in Java deniert, insbesondere ist int nicht die natural size suggested by the architecture of the execution environment1; das wäre in einer heterogenen Umgebung auch nicht machbar. double ist optional und nur benutzbar, wenn CL_DEVICE_DOUBLE_FP_CONFIG ̸= 0. Die Bits dieses Wertes zeigen die Rechenfähigkeiten von double an. OpenCL führt zusätzliche Vektortypen ein, die auf CPUs mit Vektoroperationen genutzt werden (MMX, SIMD, SSE, AVX, 3DNow!, XOP, FMA4, CVT16). Für n = 2, 3, 4, 8 und 16, sowie (u)char, (u)short, (u)int, (u)long, half, oat und double gibt es OpenCL-Typen und passende Hosttypen wie z.B. double3 und cl_double3. Mit cast- und Konstruktorschreibweise werden Variable zugewiesen:

float3 x = (float3)(3.14159265f, 2.718281828f, 0.57721566f); Mit zwei Varianten der struct-Schreibweise werden Einzelelemente aus einem Vektor herausgel"ost. Die ersten vier Elemente werden mit .x, .y, .z und .w bezeichnet. Alle (max. 16) Elemente werden mit .si, i = 0123456789abcdef indiziert. Beide Schreibweisen dürfen nicht gemischt werden:

x.y 2.71 x.z 0.577 x.zx (0.577, 3.1415) x.s0 3.1415 x.s02 (3.1415, 0.577) x.ys02 illegal Weitere Typen sind image2d_t, image3d_t, sampler_t, event_t.

13.3. Konversionen.

Implizit und explizit;

Vektortypen spärlich bis verboten. C-üblich klar, aber komplex deniert Konversionen und Reinterpretoperationen. 1

ISO/IEC9899-11,6.2.5,p.40

19

13

Zusätzliche built-ins: uchar4 u; convert_int4(u) Reelle Konversionen mit expliziten Rundungsangaben: rte, rtz, rtp, rtn (to nearest even, to zero, to plus or minus innity). float4 f; ushort4 c = convert_uchar4_sat_rte (f); reinterpret mit pointer-casts, unions. memcpy existiert nicht. auch hier: neue built-ins: float f; uint u; u = as_uint

13.4. Operatoren. 13.5. Qualier.

Vektoroperatoren

Alle Qualier mit und ohne Unterstriche.

kernel global local constant private attribute (compiler hints) access write_only read_only read_write C99: const restrict volatile

13.6. Pr"aprozessor. 13.7. Pragmas. #pragma OPENCL FP_CONTRACT ON OFF DEFAULT #pragma OPENCL EXTENSION all : enable disable

13.8. Built-ins. 14. C++ 15 16 17 18 19

Inhalt

15

Inhaltsverzeichnis Kapitel 1. Einführung in OpenCL 1. Kursinhalt und Kursbeispiele 2. Designer 3. Literatur 4. Geschichte 5. Ziele und Eigenschaften 6. Compiler und Systeme 7. OpenCL an der Universität Regensburg

3 3 3 3 4 4 5 5

Kapitel 2. crash course 1. OpenCL und C99/C++ 2. Compilation 3. C/C++-Header 4. Fehlerbehandlung 5. Abfragen 6. Prinzipieller Ablauf eines OpenCL-Programms 7. Objekte 8. Wichtige Objektbeispiele 9. OpenCL-Konzept 10. Warteschlangen 11. Events 12. Beispiel 13. Die Sprache OpenCL-C 14. C++ 15. 16. 17. 18. 19.

7 7 8 8 8 9 9 10 10 10 11 11 11 11 13 13 13 13 13 13

Inhalt

15

17

18

TTH-Seite:

http://hutchinson.belmont.ma.us/tth/

INHALTSVERZEICHNIS