Open Computing Language (OpenCL) is a framework for writing programs that execute across heterogeneous platforms consisting of central processing unit (CPUs), graphics processing unit (GPUs), 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 Intel, Advanced Micro Devices, Nvidia, and ARM Holdings.

OpenCL gives any application access to the graphics processing unit for non-graphical computing. Thus, OpenCL extends the power of the graphics processing unit beyond graphics (see general-purpose computing on graphics processing units). Academic researchers have investigated automatically compiling OpenCL programs into application-specific processors running on FPGAs,^[1] and commercial FPGA vendors are developing tools to translate OpenCL to run on their FPGA devices.^[2]

History[link]

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, Intel, and Nvidia. Apple submitted this initial proposal to the Khronos Group. On 16 June 2008, the Khronos Compute Working Group was formed^[3] 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.^[4] This technical specification was reviewed by the Khronos members and approved for public release on 8 December 2008.^[5]

OpenCL 1.0[link]

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

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 (and DirectX 11) instead of the now deprecated Close to Metal in its Stream framework.^[7][8] RapidMind announced their adoption of OpenCL underneath their development platform to support GPUs from multiple vendors with one interface.^[9] On 9 December 2008, Nvidia announced its intention to add full support for the OpenCL 1.0 specification to its GPU Computing Toolkit.^[10] On 30 October 2009, IBM released its first OpenCL implementation as a part of the XL compilers.^[11]

OpenCL 1.1[link]

OpenCL 1.1 was ratified by the Khronos Group on 14 June 2010^[12] 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[link]

On 15 November 2011, the Khronos Group announced the OpenCL 1.2 specification,^[13] 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.

Implementation[link]

• 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.^[14][15]
• 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.^[16]
• 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.^[17]
• On 20 April 2009, Nvidia announced the release of its OpenCL driver and SDK to developers participating in its OpenCL Early Access Program.^[18]
• 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.^[19]
• On 28 August 2009, Apple released Mac OS X Snow Leopard, which contains a full implementation of OpenCL.^[20]

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.^[21]

• 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.^[22]
• On 26 November 2009, Nvidia released drivers for OpenCL 1.0 (rev 48).

The Apple,^[23] Nvidia,^[24] RapidMind^[25] and Gallium3D^[26] 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.^[27]
• On 10 December 2009, VIA released their first product supporting OpenCL 1.0 - ChromotionHD 2.0 video processor included in VN1000 chipset.^[28]
• On 21 December 2009, AMD released the production version of the ATI Stream SDK 2.0,^[29] 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.^[30]
• On 30 June 2010, IBM released a fully conformant version of OpenCL 1.0.^[31]
• 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.^[32]
• 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.^[33][34]
• On 31 March 2011, IBM released a fully conformant version of OpenCL 1.1.^[31][35]
• On 25 April 2011, IBM released OpenCL Common Runtime v0.1 for Linux on x86 Architecture.^[36]
• On 4 May 2011, Nokia Research releases an open source WebCL extension for the Firefox web browser, providing a JavaScript binding to OpenCL.^[37]
• On 1 July 2011, Samsung Electronics releases an open source prototype implementation of WebCL for WebKit, providing a JavaScript binding to OpenCL.^[38]
• 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.^[39]
• On 12 December 2011, AMD released AMD APP SDK v2.6^[40] 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.^[41]
• On 17 April 2012 Khronos released a WebCL working draft^[42].

OpenCL language[link]

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.^[43] The language is extended to easily use parallelism with vector types and operations, synchronization, functions to work with work-items/groups.^[44] It has memory region qualifiers: __global, __local, __constant, and __private. Also, many built-in functions are added.

Example: computing the FFT[link]

This example will load a fast Fourier transform (FFT) implementation and execute it. The implementation is shown below. ^[45] <source lang="c">

 // 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;
 clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);

</source>

The actual calculation (based on Fitting FFT onto the G80 Architecture):^[46] <source lang="c">

 // 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);
 }

</source> A full, open source implementation of an OpenCL FFT can be found on Apple's website.^[47]

OpenCL-conformant products[link]

The Khronos Group maintains an extended list of OpenCL-conformant products.^[48]

Extensions[link]

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 SDK's. 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
• 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[link]

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[link]

The neutrality of this section is disputed. Please see the discussion on the talk page. Please do not remove this message until the dispute is resolved. (May 2012)

Multiple comparisons have been drawn between CUDA and OpenCL since its inception.^[59][60] They both draw the same conclusions: 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, unlike CUDA. CUDA is more acutely aware of the platform upon which it will be executing because it is limited to Nvidia hardware, and therefore, it provides more mature compiler optimisations and execution techniques. Furthermore, the CUDA compiler displayed more mature compilation techniques, such as more aggressive #pragma unroll addition to loops. Therefore, the developer is required to add in the optimisations manually to the kernel code. This is indicative of the maturity of the CUDA toolkit versus the newer OpenCL toolkits. It is likely in the future that this gap will be closed as the toolchains mature.

See also[link]

References[link]

External links[link]

This article's use of external links may not follow Wikipedia's policies or guidelines. Please improve this article by removing excessive or inappropriate external links, and converting useful links where appropriate into footnote references. (February 2010)

• Official website for OpenCL
• Official website for WebCL

Documentation[link]

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

Drivers[link]

• OpenCL for Nvidia (Download page)
• OpenCL for AMD (Download page)
• OpenCL for Intel (Download Page)
• OpenCL for IBM (Download page)

Libraries[link]

• OpenCLLink package for Mathematica
• OpenCL toolbox for MATLAB
• AMD Accelerated Parallel Processing Math Libraries
• ViennaCL
• GPU Automatically Tuned Linear Algebra Software
• OpenCL Data Parallel Primitives Library

Language bindings and wrappers[link]

• WebCL JavaScript bindings for Firefox
• WebCL JavaScript bindings for WebKit
• cl4d D bindings
• JOCL Java bindings
• Aparapi Java bindings
• JogAmp Java bindings
• PyOpenCL Python bindings
• PyOpenCL Python bindings
• The Open Toolkit library cross-platform C#, OpenGL, OpenAL and OpenCL wrapper for Mono/.NET
• OpenCL 1.1 headers for thinBasic
• Cloo C# bindings

Tools[link]

• gDEBugger CL debugger
• clcc Open-source, cross-platform OpenCL compiler and API entry point loader
• Fixstars OpenCL Cross Compiler Source-to-source compiler from OpenCL C to readable C code
• MinGW Package with OpenCL support
• OpenCL definition language for Notepad++
• GPU Caps Viewer
• OpenCL-Z
• Scalable Heterogeneous Computing (SHOC) Benchmark Suite
• OpenCL Studio integrated development environment
• CLyther extension language
• ImageMagick
• MandelGPU, SmallptGPU, and JuliaGPU
• Eyeon Fusion