How do you create and execute an OpenCL kernel?

To create and execute an OpenCL kernel, you generally follow these steps:

Set up OpenCL

• Find a platform and device (GPU/CPU).
• Create an OpenCL context.
• Create a command queue.

Write the kernel

OpenCL kernels are usually written in OpenCL C:

__kernel void vector_add(
    __global const float *a,
    __global const float *b,
    __global float *result)
{
    int i = get_global_id(0);
    result[i] = a[i] + b[i];
}

Create the program

Load the kernel source into your host application and create an OpenCL program:

cl_program program = clCreateProgramWithSource(
    context, 1, &kernelSource, NULL, &err);

Compile/build the kernel

err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

If compilation fails, retrieve the build log with clGetProgramBuildInfo().

Create the kernel object

cl_kernel kernel = clCreateKernel(program, "vector_add", &err);

Create and populate device buffers

cl_mem bufA = clCreateBuffer(
    context, CL_MEM_READ_ONLY,
    sizeof(float) * N, NULL, &err);

cl_mem bufB = clCreateBuffer(
    context, CL_MEM_READ_ONLY,
    sizeof(float) * N, NULL, &err);

cl_mem bufResult = clCreateBuffer(
    context, CL_MEM_WRITE_ONLY,
    sizeof(float) * N, NULL, &err);

Copy the input data to the device using clEnqueueWriteBuffer().

Set kernel arguments

clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufResult);

Execute the kernel

Specify the number of work-items and enqueue it:

size_t globalSize = N;

clEnqueueNDRangeKernel(
    queue,
    kernel,
    1,
    NULL,
    &globalSize,
    NULL,
    0,
    NULL,
    NULL);

The important distinction is that clCreateKernel() doesn't execute the kernel. It creates a kernel object. Execution happens when you call clEnqueueNDRangeKernel().

Read the results back

clEnqueueReadBuffer(
    queue,
    bufResult,
    CL_TRUE,
    0,
    sizeof(float) * N,
    result,
    0,
    NULL,
    NULL);

Overall flow

OpenCL platform
      ↓
Select device
      ↓
Create context
      ↓
Create command queue
      ↓
Load kernel source
      ↓
Create program
      ↓
Build program
      ↓
Create kernel
      ↓
Create device buffers
      ↓
Set kernel arguments
      ↓
clEnqueueNDRangeKernel()
      ↓
Read results

The host program (C/C++/Python/etc.) controls the OpenCL runtime, while the kernel runs on the selected OpenCL device.