OpenCL

This article is about the parallel computing library. For the cryptographic library, see Botan (programming library).

OpenCL

Original author(s)	Apple Inc.
Developer(s)	Khronos Group
Stable release	1.2 / 15 November 2011 (2011-11-15)
Preview release	2.0 / 22 July 2013 (2013-07-22)
Operating system	Cross-platform
Type	API
License	Unspecified
Website	www.khronos.org/opencl www.khronos.org/webcl

Open Computing Language (OpenCL) is a framework for writing programs that execute across heterogeneous platforms consisting of central processing units (CPUs), graphics processing units (GPUs), digital signal processors (DSPs) and other processors. OpenCL includes a language (based on C99) for writing kernels (functions that execute on OpenCL devices), plus application programming interfaces (APIs) that are used to define and then control the platforms. OpenCL provides parallel computing using task-based and data-based parallelism. OpenCL is an open standard maintained by the non-profit technology consortium Khronos Group. It has been adopted by Apple, Intel, Qualcomm, Advanced Micro Devices (AMD), Nvidia, Altera, Samsung, Vivante and ARM Holdings.

For example, OpenCL can be used to give an application access to a graphics processing unit for non-graphical computing (see general-purpose computing on graphics processing units). Academic researchers have investigated automatically compiling OpenCL programs into application-specific processors running on FPGAs, and commercial FPGA vendors are developing tools to translate OpenCL to run on their FPGA devices.

History

OpenCL was initially developed by Apple Inc., which holds trademark rights, and refined into an initial proposal in collaboration with technical teams at AMD, IBM, Qualcomm, Intel, and Nvidia. Apple submitted this initial proposal to the Khronos Group. On 16 June 2008, the Khronos Compute Working Group was formed with representatives from CPU, GPU, embedded-processor, and software companies. This group worked for five months to finish the technical details of the specification for OpenCL 1.0 by 18 November 2008. This technical specification was reviewed by the Khronos members and approved for public release on 8 December 2008.

OpenCL 1.0

OpenCL 1.0 has been released with Mac OS X Snow Leopard. According to an Apple press release:

Snow Leopard further extends support for modern hardware with Open Computing Language (OpenCL), which lets any application tap into the vast gigaflops of GPU computing power previously available only to graphics applications. OpenCL is based on the C programming language and has been proposed as an open standard.

AMD has decided to support OpenCL instead of the now deprecated Close to Metal in its Stream framework. RapidMind announced their adoption of OpenCL underneath their development platform to support GPUs from multiple vendors with one interface. On 9 December 2008, Nvidia announced its intention to add full support for the OpenCL 1.0 specification to its GPU Computing Toolkit. On 30 October 2009, IBM released its first OpenCL implementation as a part of the XL compilers.

OpenCL 1.1

OpenCL 1.1 was ratified by the Khronos Group on 14 June 2010 and adds significant functionality for enhanced parallel programming flexibility, functionality, and performance including:

• New data types including 3-component vectors and additional image formats;
• Handling commands from multiple host threads and processing buffers across multiple devices;
• Operations on regions of a buffer including read, write and copy of 1D, 2D, or 3D rectangular regions;
• Enhanced use of events to drive and control command execution;
• Additional OpenCL built-in C functions such as integer clamp, shuffle, and asynchronous strided copies;
• Improved OpenGL interoperability through efficient sharing of images and buffers by linking OpenCL and OpenGL events.

OpenCL 1.2

On 15 November 2011, the Khronos Group announced the OpenCL 1.2 specification, which added significant functionality over the previous versions in terms of performance and features for parallel programming. Most notable features include:

• Device partitioning: the ability to partition a device into sub-devices so that work assignments can be allocated to individual compute units. This is useful for reserving areas of the device to reduce latency for time-critical tasks.
• Separate compilation and linking of objects: the functionality to compile OpenCL into external libraries for inclusion into other programs.
• Enhanced image support: 1.2 adds support for 1D images and 1D/2D image arrays. Furthermore, the OpenGL sharing extensions now allow for OpenGL 1D textures and 1D/2D texture arrays to be used to create OpenCL images.
• Built-in kernels: custom devices that contain specific unique functionality are now integrated more closely into the OpenCL framework. Kernels can be called to use specialised or non-programmable aspects of underlying hardware. Examples include video encoding/decoding and digital signal processors.
• DirectX functionality: DX9 media surface sharing allows for efficient sharing between OpenCL and DX9 or DXVA media surfaces. Equally, for DX11, seamless sharing between OpenCL and DX11 surfaces is enabled.

OpenCL 2.0

Updates and additions to OpenCL 2.0 include (OpenCL 2.0 specification is provisional, thus still subject to change):

• Shared Virtual Memory
• Dynamic Parallelism
• Generic Address Space
• Images
• C11 Atomics
• Pipes
• Android Installable Client Driver Extension

List of vendor implementation dates

• On 10 December 2008, AMD and Nvidia held the first public OpenCL demonstration, a 75-minute presentation at Siggraph Asia 2008. AMD showed a CPU-accelerated OpenCL demo explaining the scalability of OpenCL on one or more cores while Nvidia showed a GPU-accelerated demo.
• On 16 March 2009, at the 4th Multicore Expo, Imagination Technologies announced the PowerVR SGX543MP, the first GPU of this company to feature OpenCL support.
• On 26 March 2009, at GDC 2009, AMD and Havok demonstrated the first working implementation for OpenCL accelerating Havok Cloth on AMD Radeon HD 4000 series GPU.
• On 20 April 2009, Nvidia announced the release of its OpenCL driver and SDK to developers participating in its OpenCL Early Access Program.
• On 5 August 2009, AMD unveiled the first development tools for its OpenCL platform as part of its ATI Stream SDK v2.0 Beta Program.
• On 28 August 2009, Apple released Mac OS X Snow Leopard, which contains a full implementation of OpenCL.

OpenCL in Snow Leopard is supported on the Nvidia GeForce 320M, GeForce GT 330M, GeForce 9400M, GeForce 9600M GT, GeForce 8600M GT, GeForce GT 120, GeForce GT 130, GeForce GTX 285, GeForce 8800 GT, GeForce 8800 GS, Quadro FX 4800, Quadro FX5600, ATI Radeon HD 4670, ATI Radeon HD 4850, Radeon HD 4870, ATI Radeon HD 5670, ATI Radeon HD 5750, ATI Radeon HD 5770 and ATI Radeon HD 5870.

• On 28 September 2009, Nvidia released its own OpenCL drivers and SDK implementation.
• On 13 October 2009, AMD released the fourth beta of the ATI Stream SDK 2.0, which provides a complete OpenCL implementation on both R700/R800 GPUs and SSE3 capable CPUs. The SDK is available for both Linux and Windows.
• On 26 November 2009, Nvidia released drivers for OpenCL 1.0 (rev 48).

The Apple, Nvidia, RapidMind and Gallium3D implementations of OpenCL are all based on the LLVM Compiler technology and use the Clang Compiler as its frontend.

• On 27 October 2009, S3 released their first product supporting native OpenCL 1.0 - the Chrome 5400E embedded graphics processor.
• On 10 December 2009, VIA released their first product supporting OpenCL 1.0 - ChromotionHD 2.0 video processor included in VN1000 chipset.
• On 21 December 2009, AMD released the production version of the ATI Stream SDK 2.0, which provides OpenCL 1.0 support for R800 GPUs and beta support for R700 GPUs.
• On 1 June 2010, ZiiLABS released details of their first OpenCL implementation for the ZMS processor for handheld, embedded and digital home products.
• On 30 June 2010, IBM released a fully conformant version of OpenCL 1.0.
• On 13 September 2010, Intel released details of their first OpenCL implementation for the Sandy Bridge chip architecture. Sandy Bridge will integrate Intel's newest graphics chip technology directly onto the central processing unit.
• On 15 November 2010, Wolfram Research released Mathematica 8 with OpenCLLink package.
• On 3 March 2011, Khronos Group announces the formation of the WebCL working group to explore defining a JavaScript binding to OpenCL. This creates the potential to harness GPU and multi-core CPU parallel processing from a Web browser.
• On 31 March 2011, IBM released a fully conformant version of OpenCL 1.1.
• On 25 April 2011, IBM released OpenCL Common Runtime v0.1 for Linux on x86 Architecture.
• On 4 May 2011, Nokia Research releases an open source WebCL extension for the Firefox web browser, providing a JavaScript binding to OpenCL.
• On 1 July 2011, Samsung Electronics releases an open source prototype implementation of WebCL for WebKit, providing a JavaScript binding to OpenCL.
• On 8 August 2011, AMD released the OpenCL-driven AMD Accelerated Parallel Processing (APP) Software Development Kit (SDK) v2.5, replacing the ATI Stream SDK as technology and concept.
• On 12 December 2011, AMD released AMD APP SDK v2.6 which contains a preview of OpenCL 1.2.
• On 27 February 2012, The Portland Group released the PGI OpenCL compiler for multi-core ARM CPUs.
• On 17 April 2012 Khronos released a WebCL working draft.

OpenCL language

The programming language used to write computation kernels is based on C99 with some limitations and additions. It omits the use of function pointers, recursion, bit fields, variable-length arrays, and standard C99 header files. The language is extended to easily use parallelism with vector types and operations, synchronization, functions to work with work-items/groups. It has memory region qualifiers: __global, __local, __constant, and __private. Also, many built-in functions are added.

Example: computing the FFT

This example will load a fast Fourier transform (FFT) implementation and execute it. The implementation is shown below.

  // create a compute context with GPU device
  context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
 
  // create a command queue
  clGetDeviceIDs( NULL, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL );
  queue = clCreateCommandQueue(context, device_id, 0, NULL);
 
  // allocate the buffer memory objects
  memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA, NULL);
  memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL, NULL);
 
  // create the compute program
  program = clCreateProgramWithSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL);
 
  // build the compute program executable
  clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
 
  // create the compute kernel
  kernel = clCreateKernel(program, "fft1D_1024", NULL);
 
  // set the args values
  clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs[0]);
  clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs[1]);
  clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size[0]+1)*16, NULL);
  clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size[0]+1)*16, NULL);
 
  // create N-D range object with work-item dimensions and execute kernel
  global_work_size[0] = num_entries;
  local_work_size[0] = 64; //Nvidia: 192 or 256
  clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);

The actual calculation (based on Fitting FFT onto the G80 Architecture):

  // This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into
  // calls to a radix 16 function, another radix 16 function and then a radix 4 function
 
  __kernel void fft1D_1024 (__global float2 *in, __global float2 *out,
                          __local float *sMemx, __local float *sMemy) {
    int tid = get_local_id(0);
    int blockIdx = get_group_id(0) * 1024 + tid;
    float2 data[16];
 
    // starting index of data to/from global memory
    in = in + blockIdx;  out = out + blockIdx;
 
    globalLoads(data, in, 64); // coalesced global reads
    fftRadix16Pass(data);      // in-place radix-16 pass
    twiddleFactorMul(data, tid, 1024, 0);
 
    // local shuffle using local memory
    localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4)));
    fftRadix16Pass(data);               // in-place radix-16 pass
    twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication
 
    localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15)));
 
    // four radix-4 function calls
    fftRadix4Pass(data);      // radix-4 function number 1
    fftRadix4Pass(data + 4);  // radix-4 function number 2
    fftRadix4Pass(data + 8);  // radix-4 function number 3
    fftRadix4Pass(data + 12); // radix-4 function number 4
 
    // coalesced global writes
    globalStores(data, out, 64);
  }

A full, open source implementation of an OpenCL FFT can be found on Apple's website.

OpenCL-conformant products

The Khronos Group maintains an extended list of OpenCL-conformant products.

Synopsis of OpenCL conformant products
AMD APP SDK (supports OpenCL CPU and accelerated processing unit Devices)	X86 + SSE2 (or higher) compatible CPUs 64-bit & 32-bit; Linux 2.6 PC, Windows Vista/7 PC	AMD Fusion E-350, E-240, C-50, C-30 with HD 6310/HD 6250	AMD Radeon/Mobility HD 6800, HD 5x00 series GPU, iGPU HD 6310/HD 6250	ATI FirePro Vx800 series GPU
Intel SDK for OpenCL Applications 2012 (supports Intel Core processors and Intel HD Graphics 4000/2500)	Intel CPUs with SSE 4.1, SSE 4.2 or AVX support. Microsoft Windows, Linux	Intel Core i7, i5, i3; 2nd Generation Intel Core i7/5/3, 3rd Generation Intel Core Processors with Intel HD Graphics 4000/2500	Intel Core 2 Solo, Duo Quad, Extreme	Intel Xeon 7x00,5x00,3x00 (Core based)
IBM Servers with OpenCL Development Kit for Linux on Power running on Power VSX	IBM Power 755 (PERCS), 750	IBM BladeCenter PS70x Express	IBM BladeCenter JS2x, JS43	IBM BladeCenter QS22
IBM OpenCL Common Runtime (OCR)	X86 + SSE2 (or higher) compatible CPUs 64-bit & 32-bit; Linux 2.6 PC	AMD Fusion, Nvidia Ion and Intel Core i7, i5, i3; 2nd Generation Intel Core i7/5/3	AMD Radeon, Nvidia GeForce and Intel Core 2 Solo, Duo, Quad, Extreme	ATI FirePro, Nvidia Quadro and Intel Xeon 7x00,5x00,3x00 (Core based)
Nvidia OpenCL Driver and Tools	Nvidia Tesla C/D/S	Nvidia GeForce GTS/GT/GTX	Nvidia Ion	Nvidia Quadro FX/NVX/Plex

Extensions

Some vendors provide extended functionality over the standard OpenCL specification via the means of extensions. These are still specified by Khronos but provided by vendors within their SDKs. They often contain features that are to be implemented in the future - for example device fission functionality was originally an extension but is now provided as part of the 1.2 specification.

Extensions provided in the 1.2 specification include:

• Writing to 3D image memory objects
• Half-precision floating-point format
• Sharing memory objects with OpenGL
• Creating event objects from GL sync objects
• Sharing memory objects with Direct3D 10
• DX9 media Surface Sharing
• Sharing Memory Objects with Direct3D 11

Device fission

Device fission - introduced fully into the OpenCL standard with version 1.2 - allows individual command queues to be used for specific areas of a device. For example, within the Intel SDK, a command queue can be created that maps directly to an individual core. AMD also provides functionality for device fission, also originally as an extension. Device fission can be used where the availability of compute is required reliably, such as in a latency sensitive environment. Fission effectively reserves areas of the device for computation.

Comparison with CUDA

Two comparisons have been drawn between CUDA and OpenCL since its inception, and both arrived at the same conclusion: if the OpenCL implementation is correctly tweaked to suit the target architecture, it performs no worse than CUDA. Because the key feature of OpenCL is portability (via its abstracted memory and execution model), the programmer is not able to directly use GPU-specific technologies (such as inline PTX) unless they are willing to give up direct portability. CUDA is more directly connected to the platform upon which it will be executing because it is limited to Nvidia hardware. Compiler technology for both standards supported by Nvidia's toolkit is based upon LLVM and compiles to Nvidia's PTX instruction set abstraction.

See also

• GPGPU
• OpenGL
• CUDA
• OpenHMPP
• OpenACC
• DirectCompute
• FireStream
• Larrabee
• Close to Metal
• BrookGPU
• Lib Sh
• CLyther
• SIMD
• C++ AMP
• SequenceL

External links

• Official website for OpenCL
• Official website for WebCL

Documentation

• OpenCL 1.2 Specification
• OpenCL 1.2 Reference Pages
• OpenCL 1.2 Quick Reference card

Drivers

• OpenCL for Nvidia (Download page)
• OpenCL for AMD (Download page)
• OpenCL for Intel (Download Page)
• OpenCL for IBM (Download page)
• pocl, a generic open source implementation based on Clang and LLVM (Download page)

Libraries

• OpenCLLink: package for Mathematica
• OpenCL at Google Code: toolbox for MATLAB
• AMD Accelerated Parallel Processing Math Libraries
• ViennaCL on SourceForge.net
• VexCL: C++ vector expression template library for OpenCL
• GPU Automatically Tuned Linear Algebra Software
• OpenCL Data Parallel Primitives Library at Google Code
• Asgard Abstraction layer to simplify OpenCL
• Boost.Compute: C++ wrapper and STL-like API for OpenCL

Language bindings and wrappers

JavaScript - WebCL:

• node-webcl - JavaScript bindings for Node.js
• Nokia WebCL extension: JavaScript bindings for Firefox
• WebCL for WebKit at Google Code: JavaScript bindings for WebKit

Java:

• Aparapi at Google Code: Java bindings
• JogAmp (JOCL) : Java bindings
• LWJGL

Python:

• PyOpenCL
• CLyther

Ruby:

• No bindings are actually maintained[1]

Rust:

• rust-opencl : Rust bindings

D:

• cl4d: D bindings

C#:

• The Open Toolkit library on SourceForge.net: cross-platform C#, OpenGL, OpenAL and OpenCL wrapper for Mono/.NET
• Cloo on SourceForge.net: C# bindings

thinBasic:

• OpenCL 1.1 headers for thinBasic

Tools

• AMD CodeXL - OpenCL debugger and profiler
• clcc on SourceForge.net: Open-source, cross-platform OpenCL compiler and API entry point loader
• CAPS Compilers: An OpenCL code generator out of OpenACC and OpenHMPP directives in C/C++ or Fortran
• OpenCL definition language for Notepad++
• GPU Caps Viewer
• OpenCL-Z on SourceForge.net
• Scalable Heterogeneous Computing (SHOC) Benchmark Suite
• OpenCL Studio: Integrated development environment
• CLyther on SourceForge.net: extension language
• ImageMagick
• MandelGPU, SmallptGPU, and JuliaGPU
• Eyeon Fusion
• OpenCL CodeBench: host code generator and OpenCL-aware kernel editor for Eclipse
• VirtualCL (VCL) cluster platform: a wrapper for OpenCL that allows most unmodified applications to transparently utilize multiple OpenCL devices in a cluster as if all the devices are on the local computer.