OpenCL

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.

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. // 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 The kernel function computes, in each invocation, the dot product of a single row of a matrix and a vector : y_i = a_{i,:} \cdot x = \sum_j a_{i,j} x_j . 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 . include include 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[1]; 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], 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[1] = { }; 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[1] = { 256 }; 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 size_t global_work_size[1] = { 256 }; 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 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[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 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. The following code snippet illustrates how kernels with complex-number arithmetic can be implemented in C++ for OpenCL language with convenient use of C++ features. // Define a class Complex, that can perform complex-number computations with // various precision when different types for T are used - double, float, half. template class complex_t { T m_re; // Real component. T m_im; // Imaginary component. public: complex_t(T re, T im): m_re{re}, m_im{im} {}; // Define operator for complex-number multiplication. complex_t operator*(const complex_t &other) const { return {m_re * other.m_re - m_im * other.m_im, m_re * other.m_im + m_im * other.m_re}; } T get_re() const { return m_re; } T get_im() const { return m_im; } }; // A helper function to compute multiplication over complex numbers read from // the input buffer and to store the computed result into the output buffer. template void compute_helper(__global T *in, __global T *out) { auto idx = get_global_id(0); // Every work-item uses 4 consecutive items from the input buffer // - two for each complex number. auto offset = idx * 4; auto num1 = complex_t{in[offset], in[offset + 1]}; auto num2 = complex_t{in[offset + 2], in[offset + 3]}; // Perform complex-number multiplication. auto res = num1 * num2; // Every work-item writes 2 consecutive items to the output buffer. out[idx * 2] = res.get_re(); out[idx * 2 + 1] = res.get_im(); } // This kernel is used for complex-number multiplication in single precision. __kernel void compute_sp(__global float *in, __global float *out) { compute_helper(in, out); } ifdef cl_khr_fp16 // This kernel is used for complex-number multiplication in half precision when // it is supported by the device. pragma OPENCL EXTENSION cl_khr_fp16: enable __kernel void compute_hp(__global half *in, __global half *out) { compute_helper(in, out); } endif C++ for OpenCL sources can be compiled by OpenCL drivers that support cl_ext_cxx_for_opencl extension, which allows the use of -cl-std=CLC++ in clBuildProgram() . Arm has announced support for this extension in December 2020. However, due to increasing complexity of the algorithms accelerated on OpenCL devices, it is expected that more applications will compile C++ for OpenCL kernels offline using stand alone compilers such as Clang into executable binary format or portable binary format e.g. SPIR-V. Such an executable can be loaded during the OpenCL applications execution using a dedicated OpenCL API. Binaries compiled from sources in C++ for OpenCL 1.0 can be executed on OpenCL 2.0 conformant devices. Depending on the language features used in such kernel sources it can also be executed on devices supporting earlier OpenCL versions or OpenCL 3.0. Aside from OpenCL drivers, kernels written in C++ for OpenCL can be compiled for execution on Vulkan devices using clspv compiler and clvk runtime layer, just the same way as OpenCL C kernels.

Text derived from https://en.wikipedia.org/w/index.php?oldid=1378554601 under CC BY-SA 4.0; modifications have been made.