OpenCL Objects, Contexts, and Runtime Compilation

OpenCL was designed from the outset as a compute interface spanning devices from more than one vendor, and its object model reads like the checklist of everything that had to be made explicit to achieve that. Where a single-vendor API can fold context creation into the first allocation, OpenCL separates it, because the platform beneath it may be a discrete GPU, an integrated GPU sharing host memory, a CPU with vector units, or a DSP. The cost of that explicitness is a long call sequence before the first work-item runs.

The object graph is the API

Five object types, layered:

A platform is a vendor’s implementation of the runtime. clGetPlatformIDs enumerates them, and a machine with an Intel CPU, an AMD GPU, and an NVIDIA GPU may present three. A platform is the unit of vendor identity, not the unit of hardware.

A device belongs to a platform. clGetDeviceIDs enumerates them and takes a device type (CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_ACCELERATOR, or CL_DEVICE_TYPE_DEFAULT), which is the API’s first hint that the CPU is a legitimate target with its own implementation of the same model.

A context binds a set of devices from one platform, and everything else is created within it. Buffers, programs, kernels, and queues all carry a reference to their context, and an object cannot be used with an object from another context. That rule is a hard error, not a soft one, and it is the reason a library that creates its own context cannot share buffers with its caller.

A command queue is attached to a single device within the context. This is where the context’s multi-device capability stops being useful: work is enqueued to one device at a time, and distributing a kernel across two devices means two queues and explicit partitioning of the data, not one launch.

Programs and kernels are the last layer, and they are the part with the surprising structure. The four obligations the whole object graph serves are set out in The Contract a Compute API Encodes.

A program is source plus one or more binaries

clCreateProgramWithSource takes OpenCL C source as an array of strings with a parallel array of lengths, so a program can be assembled from several source strings without concatenating them first. clBuildProgram then compiles it for a list of devices in the context: possibly more than one device, possibly devices with different instruction sets, from one build call.

The build produces a program object holding one binary per device. The program is not a kernel; clCreateKernel(program, "name", &err) looks up an entry point by name in the program’s exported symbol table and produces a kernel object that can be given arguments. Two kernels from the same program share its binaries and its compiled constant data but have independent argument states.

The build is where portability is decided, and it is decided at run time on the user’s machine. clBuildProgram accepts an options string that is passed to the device compiler, and its contents are vendor-specific: which -cl-* flags are honoured, what -D macros are defined by default, what the default optimisation level is, whether -cl-fast-relaxed-math changes results. A build that succeeds under one vendor’s compiler may fail under another’s, and the error surfaces at run time with a build log rather than at development time.

/* The count is a kernel argument, so the guard tests the caller's real size
   even when the launch rounds the global size up to a multiple of the local size. */
static const char *src =
    "__kernel void axpy(__global float *y, __global const float *x,\n"
    "                   const float a, const int n) {\n"
    "    int i = (int)get_global_id(0);\n"
    "    if (i < n) y[i] += a * x[i];\n"
    "}\n";

cl_int err;
cl_program prog = clCreateProgramWithSource(ctx, 1, &src, NULL, &err);
err = clBuildProgram(prog, 1, &dev, "-cl-mad-enable", NULL, NULL);
if (err != CL_SUCCESS) {
    size_t len = 0;
    clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &len);
    char *log = malloc(len);
    clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, len, log, NULL);
    fputs(log, stderr);            /* the only place the reason appears */
    free(log);
}

The two-call pattern (query the length, allocate, query the contents) recurs for every variable-length property in OpenCL, from the build log to the device name to the list of supported extensions. It is the single most visible ergonomic cost of a C API designed around opaque handles, and it applies to the compiler diagnostics, which means a build failure without a log dump leaves the program with no information about why.

Kernel arguments are copied, and the copy is typed

clSetKernelArg(kernel, index, sizeof(arg), &arg) takes a size and a pointer. For a scalar the size is the scalar’s size. For a buffer the size is sizeof(cl_mem), the size of the handle rather than of the buffer, and the pointer is to the handle. This asymmetry catches out every new caller once, and the API cannot detect it: passing the buffer’s byte size where the handle size was expected is a plausible-looking integer, and the runtime either rejects the argument count later or reads the wrong number of bytes.

Local memory arguments are the third case. Passing NULL as the value with a size of the desired byte count declares a dynamic local array, sized at launch rather than at compile time:

clSetKernelArg(k, 0, sizeof(cl_mem), &input);
clSetKernelArg(k, 1, 64 * sizeof(float), NULL);  // dynamic __local tile

The size must not exceed CL_DEVICE_LOCAL_MEM_SIZE, and the sum of all local arguments plus any statically declared __local arrays must fit the per-work-group budget, which is usually smaller than the per-device total. The distinction matters because a work-group’s local usage gates how many groups can be resident.

The NDRange call and its three shapes

clEnqueueNDRangeKernel takes a work dimension from 1 to 3, a global size array, an optional local size array, and an optional global offset.

When local size is NULL the implementation chooses, which removes a portability burden but also removes control over the two properties that depend on it: work-group-local synchronisation and local memory budgeting. When it is supplied, OpenCL 1.x requires the global size to be a multiple of the local size and returns CL_INVALID_WORK_GROUP_SIZE otherwise; later versions relax that only where the device advertises support for non-uniform work-groups. The guard predicate idiom from CUDA therefore applies to portable code rather than being merely advisable. The remedy is to round the global size up to a multiple of the local size and guard inside the kernel, which is why the sample above takes the element count as an argument: the rounded launch reports the rounded size through get_global_size, so the guard must compare against the caller’s real count and not against the launch extent.

The global offset shifts the origin of the index space without changing its extent. It exists so that a kernel can be pointed at a sub-range of an array without adjusting its arithmetic, and it is the only place in the API where the index space is not assumed to start at zero.

Where the portability actually leaks

Three specific places, all of them visible once the object graph is understood.

The first is the compiler. OpenCL C is a subset of C99 with added vector types and address-space qualifiers, and the address-space qualifiers are mandatory on pointer parameters; implementations differ in how strictly they enforce that. The specification’s build options also differ in support: -cl-std selects the language version, and implementations vary in which versions they accept.

The second is floating-point contraction. Whether a * b + c is contracted into a fused multiply-add is implementation-defined, and -cl-mad-enable is the build option that permits it explicitly; -cl-fp32-correctly-rounded-divide-sqrt is a separate guarantee about division and square root. The results of a reduction can differ in the last bits across devices as a consequence. Code that compares floating-point results for equality across two platforms will fail.

The third is the device itself. clGetDeviceInfo reports a long list of properties, among them CL_DEVICE_MAX_COMPUTE_UNITS, CL_DEVICE_MAX_WORK_GROUP_SIZE, CL_DEVICE_GLOBAL_MEM_SIZE, and CL_DEVICE_LOCAL_MEM_SIZE, and the portable program queries all of them rather than assuming. On a CPU implementation the compute units are cores and the local memory is an ordinary cache, so a kernel tuned to use local memory as scratch may simply not benefit; the model is uniform, the hardware is not.

Programs can be cached, and caching is the practical answer

Because the build happens on the target machine and its cost lands on first use, clGetProgramInfo with CL_PROGRAM_BINARIES retrieves the compiled binary per device, and clCreateProgramWithBinary reconstructs a program from it. The binary is only valid for the device it was built for. The device’s CL_DEVICE_VERSION and CL_DEVICE_NAME should be recorded alongside it, and the specification allows an implementation to reject a binary it does not recognise.

This binary cache is the pattern behind the dispatch layers described in Capability Negotiation Across Architectures: probe the device once, compile or load a binary keyed to it, cache the result, and keep the source as the fallback that always works. The runtime compilation model is a liability until it is treated as a cache, at which point it becomes the mechanism that lets one binary distribution serve several generations of hardware.