OpenCL (Open Computing Language) 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), field-programmable gate arrays (FPGAs) and other processors or hardware accelerators. OpenCL specifies a programming language (based on C99) for programming these devices and application programming interfaces (APIs) to control the platform and execute programs on the compute devices. OpenCL provides a standard interface for parallel computing using task- and data-based parallelism.
OpenCL is an open standard maintained by the Khronos Group, a non-profit, open standards organization. Conformant implementations (passed the Conformance Test Suite) are available from a range of companies including AMD, Arm, Cadence, Google, Imagination, Intel, Nvidia, Qualcomm, Samsung, SPI and Verisilicon.
Overview OpenCL views a computing system as consisting of a number of compute devices, which might be central processing units (CPUs) or "accelerators" such as graphics processing units (GPUs), attached to a host processor (a CPU). It defines a C-like language for writing programs. Functions executed on an OpenCL device are called "kernels". A single compute device typically consists of several compute units, which in turn comprise multiple processing elements (PEs). A single kernel execution can run on all or many of the PEs in parallel. How a compute device is subdivided into compute units and PEs is up to the vendor; a compute unit can be thought of as a "core", but the notion of core is hard to define across all the types of devices supported by OpenCL (or even within the category of "CPUs"), and the number of compute units may not correspond to the number of cores claimed in vendors' marketing literature (which may actually be counting SIMD lanes).
In addition to its C-like programming language, OpenCL defines an application programming interface (API) that allows programs running on the host to launch kernels on the compute devices and manage device memory, which is (at least conceptually) separate from host memory. Programs in the OpenCL language are intended to be compiled at run-time, so that OpenCL-using applications are portable between implementations for various host devices. The OpenCL standard defines host APIs for C and C++; third-party APIs exist for other programming languages and platforms such as Python, Java, Perl, D and .NET. An implementation of the OpenCL standard consists of a library that implements the API for C and C++, and an OpenCL C compiler for the compute devices targeted.
In order to open the OpenCL programming model to other languages or to protect the kernel source from inspection, the Standard Portable Intermediate Representation (SPIR) can be used as a target-independent way to ship kernels between a front-end compiler and the OpenCL back-end.
More recently Khronos Group has ratified SYCL, a higher-level programming model for OpenCL as a single-source eDSL based on pure C++17 to improve programming productivity. People interested by C++ kernels but not in the SYCL single-source programming style can use C++ features with compute kernel sources written in "C++ for OpenCL" language.
Memory hierarchy OpenCL defines a four-level memory hierarchy for the compute device:
global memory: shared by all processing elements, but has high access latency (); read-only memory: smaller, low latency, writable by the host CPU but not the compute devices (); local memory: shared by a group of processing elements (); per-element private memory (registers; ).
Not every device needs to implement each level of this hierarchy in hardware. Consistency between the various levels in the hierarchy is relaxed, and only enforced by explicit synchronization constructs, notably barriers.
Devices may or may not share memory with the host CPU. The host API provides handles on device memory buffers and functions to transfer data back and forth between host and devices.
OpenCL kernel language The programming language that is used to write compute kernels is called kernel language. OpenCL adopts C/C++-based languages to specify the kernel computations performed on the device with some restrictions and additions to facilitate efficient mapping to the heterogeneous hardware resources of accelerators. Traditionally OpenCL C was used to program the accelerators in OpenCL standard, later C++ for OpenCL kernel language was developed that inherited all functionality from OpenCL C but allowed to use C++ features in the kernel sources.
OpenCL C language OpenCL C is a C99-based language dialect adapted to fit the device model in OpenCL. Memory buffers reside in specific levels of the memory hierarchy, and pointers are annotated with the region qualifiers , , , and , reflecting this. Instead of a device program having a function, OpenCL C functions are marked to signal that they are entry points into the program to be called from the host program. Function pointers, bit fields and variable-length arrays are omitted, and recursion is forbidden. The C standard library is replaced by a custom set of standard functions, geared toward math programming.
OpenCL C is extended to facilitate use of parallelism with vector types and operations, synchronization, and functions to work with work-items and work-groups. In particular, besides scalar types such as and , which behave similarly to the corresponding types in C, OpenCL provides fixed-length vector types such as (4-vector of single-precision floats); such vector types are available in lengths two, three, four, eight and sixteen for various base types. Vectorized operations on these types are intended to map onto SIMD instructions sets, e.g., SSE or VMX, when running OpenCL programs on CPUs. Other specialized types include 2-d and 3-d image types.
Example: matrix–vector multiplication
The following is a matrix–vector multiplication algorithm in OpenCL C.
// Multiplies A*x, leaving the result in y. // A is a row-major matrix, meaning the (i,j) element is at A[i*ncols+j]. __kernel void matvec(__global const float *A, __global const float *x, uint ncols, __global float *y) { size_t i = get_global_id(0); // Global id, used as the row index __global float const *a = &A[i*ncols]; // Pointer to the i'th row float sum = 0.f; // Accumulator for dot product for (size_t j = 0; j < ncols; j++) { sum += a[j] * x[j];
y[i] = sum; }
The kernel function computes, in each invocation, the dot product of a single row of a matrix and a vector :
To extend this into a full matrix–vector multiplication, the OpenCL runtime maps the kernel over the rows of the matrix. On the host side, the function does this; it takes as arguments the kernel to execute, its arguments, and a number of work-items, corresponding to the number of rows in the matrix .
Example: computing the FFT This example will load a fast Fourier transform (FFT) implementation and execute it. The implementation is shown below. The code asks the OpenCL library for the first available graphics card, creates memory buffers for reading and writing (from the perspective of the graphics card), JIT-compiles the FFT-kernel and then finally asynchronously runs the kernel. The result from the transform is not read in this example. This is illustrative example code, not intended for serious use, so error-handling is entirely omitted.
#include <stdio.h> #include <time.h> #include "CL/opencl.h"
#define NUM_ENTRIES 1024
int main() // (int argc, const char* argv[]) { // CONSTANTS // The source code of the kernel is represented as a string // located inside file: "fft1D_1024_kernel_src.cl". For the details see the next listing. const char *KernelSource = #include "fft1D_1024_kernel_src.cl"
// Looking up the available GPUs const cl_uint num = 1; clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 0, NULL, (cl_uint*)&num);
cl_device_id devices; clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, num, devices, NULL);
// create a compute context with GPU device cl_context context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
// create a command queue clGetDeviceIDs(NULL, CL_DEVICE_TYPE_DEFAULT, 1, devices, NULL); cl_command_queue queue = clCreateCommandQueue(context, devices, 0, NULL);
// allocate the buffer memory objects cl_mem memobjs[] = { clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * 2 * NUM_ENTRIES, NULL, NULL), clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * 2 * NUM_ENTRIES, NULL, NULL) };
// create the compute program // const char* fft1D_1024_kernel_src = { }; cl_program program = clCreateProgramWithSource(context, 1, (const char **)& KernelSource, NULL, NULL);
// build the compute program executable clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
// create the compute kernel cl_kernel kernel = clCreateKernel(program, "fft1D_1024", NULL);
// set the args values
size_t local_work_size = { 256 };
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs); clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs); clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size + 1) * 16, NULL); clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size + 1) * 16, NULL);
// create N-D range object with work-item dimensions and execute kernel size_t global_work_size = { 256 };
global_work_size = NUM_ENTRIES; local_work_size = 64; //Nvidia: 192 or 256 clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); }
The actual calculation inside file "fft1D_1024_kernel_src.cl" (based on "Fitting FFT onto the G80 Architecture"): R"( // 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;
// 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 C++ language
OpenCL C++ is a short-lived specification for a language that combines OpenCL C and C++14. It was intended to be built in an online mode only, by passing the -cl-std=c++ parameter in clBuildProgram(). No extension for detecting support for this language is described. It is unknown whether any driver actually supported this language.
C++ for OpenCL language