In May 2026, on the eve of IWOCL in Heilbronn, the Khronos OpenCL Working Group released OpenCL 3.1. It is the first version bump in almost six years, and the first since 2017 that makes the specification bigger rather than more permissive.
Two months later the first conformant 3.1 implementation appeared on the Khronos list: Apple M1 and M2 graphics, running Linux, driven by Mesa’s Rusticl.
Not Apple’s driver: Apple deprecated OpenCL in 2018 and never looked back. A reverse-engineered kernel driver plus a Rust userspace led by one engineer at Red Hat.
That is a reasonable summary of what OpenCL has become. It is not the thing that lost to CUDA and quietly died. It is a specification with a governance process that keeps grinding forward, implemented today mostly by people who are not the vendors that originally wrote it, running the mobile inference stacks of two of the largest silicon companies on earth, and serving as the portability substrate underneath SYCL, chipStar, and a lot of code that will never mention OpenCL in its README.
There is one constraint that explains almost everything about OpenCL, and it is worth putting up front because the rest of this article is evidence for it.
OpenCL is the only compute API that had to describe hardware it could not see. CUDA describes one company’s roadmap, and that company knows what silicon is coming. Vulkan compute describes graphics hardware. Metal describes Apple’s.
OpenCL signed a contract covering CPUs, DSPs, FPGAs, GPUs from four vendors, and processors that did not exist when the contract was written. The price of that contract is that every guarantee has to be weak enough for the weakest member, or optional.
Once you see that, the history stops looking like a series of unforced errors. OpenCL 2.0 tried to make strong guarantees universal, and half the industry declined to implement them. OpenCL 3.0 stopped pretending and made almost everything optional, which read as capitulation and was actually an accurate description of what already existed.
OpenCL 3.1, in May 2026, is the first release in the standard’s history that mostly ratifies what had already shipped rather than predicting what ought to.
The performance-portability problem is the same constraint wearing different clothes: a kernel is a description of a memory hierarchy, and OpenCL cannot tell you which one you have.
I’ve tried to make this concrete rather than rhetorical. Where the article makes a claim about behaviour, I actually ran it: the kernels here were compiled and executed, the accuracy bounds were measured against what an implementation actually delivers, and the compile-time and dispatch-overhead arguments have numbers attached.
Those measurements come from one CPU device and are labelled as such: they demonstrate mechanisms, not hardware.
I’ve also been fussy about versions and dates, because a large fraction of what is written about OpenCL online describes the world of 2015 and says so nowhere.
The origins
The immediate ancestors of OpenCL are worth naming because they explain its shape.
Between 2006 and 2008 there were at least four serious attempts to make GPUs programmable for non-graphics work: Stanford’s BrookGPU and its commercial descendant at AMD, ATI’s Close to Metal, NVIDIA’s CUDA (released February 2007), and IBM’s toolchain for the Cell Broadband Engine. Each was tied to one vendor’s hardware.
Apple, which at the time shipped machines with GPUs from both NVIDIA and ATI and had a strategic interest in not being locked to either, wanted one API that worked on both plus the CPU.
Apple developed an initial proposal internally, with Aaftab Munshi, who had edited the OpenGL ES 1.1 and 2.0 specifications, as spec editor, and refined it with technical teams at AMD, IBM, Qualcomm, Intel, and NVIDIA before submitting it to Khronos.
The Khronos Compute Working Group was formed on 16 June 2008. It finished the technical content of OpenCL 1.0 on 18 November 2008, and the specification was approved for public release on 8 December 2008.
Five months from working group formation to a ratified cross-vendor standard is fast by any standards-body measure, and it shows in the result: OpenCL 1.0 has the feel of a design that was mostly finished before the committee got involved.
Apple demonstrated a beta at WWDC in June 2008, a grid of 64 emulated Apple II machines running on the CPU to show task parallelism, and an N-body simulation on a Mac Pro’s GPU to show data parallelism, and shipped the first production implementation in Mac OS X 10.6 Snow Leopard on 28 August 2009.
AMD abandoned Close to Metal and backed OpenCL. NVIDIA announced OpenCL support for its GPU Computing Toolkit the day after ratification and shipped drivers in September 2009. IBM shipped an implementation for POWER through its XL compilers in October 2009. For about eighteen months, OpenCL looked like it was going to be the way GPUs got programmed.
The 1.x line then filled in the obvious gaps:
OpenCL 1.1 (14 June 2010) added three-component vector types, sub-buffers, rectangular region read/write/copy for buffers, user events, async_work_group_strided_copy, and tighter OpenGL interop through event linking. It also allowed API calls from multiple host threads, which 1.0 had not required.
OpenCL 1.2 (15 November 2011) is the version that still matters most, because it became the mandatory baseline of OpenCL 3.0 and is therefore the safe target for portable code even today.
It added device partitioning (splitting a device into sub-devices along compute-unit or cache-hierarchy boundaries), separate compilation and linking of programs, 1D images and 1D/2D image arrays, built-in kernels and custom devices, clEnqueueMigrateMemObjects, clEnqueueFillBuffer/clEnqueueFillImage, DirectX 9 media surface and DirectX 11 sharing, and -cl-fp32-correctly-rounded-divide-sqrt for code that needs IEEE 754 semantics on single-precision division and square root rather than the looser default.
By the end of 1.2 the API had a coherent shape. What happened next is the interesting part, and we’ll come back to it.
The four models
The specification organises itself around four models, and it is worth internalising them in the order the spec gives them, because most confusion about OpenCL comes from conflating two of them.
Platform model
A host is connected to one or more compute devices. Each device has one or more compute units, each of which has one or more processing elements. That is the entire hardware abstraction, and it is deliberately vague.
The mapping to real silicon is vendor-defined and it’s not always intuitive: on an NVIDIA GPU a compute unit is a streaming multiprocessor; on AMD GCN it is a compute unit containing four 16-wide SIMDs; on a CPU it is typically a hardware thread; on Adreno it is a shader processor.
CL_DEVICE_MAX_COMPUTE_UNITS therefore does not mean the same thing across vendors, and comparing it across vendors is meaningless. The number vendors quote in marketing material, “5888 cores”, usually counts processing elements or SIMD lanes, not compute units.
Devices are typed: CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_CUSTOM (1.2+), plus CL_DEVICE_TYPE_DEFAULT and CL_DEVICE_TYPE_ALL.
Custom devices are the interesting outlier: they are devices that do not support the OpenCL C programming language at all and expose only built-in kernels, fixed-function or firmware-defined entry points queried through CL_DEVICE_BUILT_IN_KERNELS. This is how OpenCL accommodates video encoders, ISPs, and fixed-function DSP blocks.
There is a trap here worth naming, because it caught me. Through OpenCL 1.2 and the 2.x line the specification defined CL_DEVICE_TYPE_ALL as every device except custom ones, and that is still what most tutorials say.
The current unified specification does not: CL_DEVICE_TYPE_ALL is now simply “all OpenCL devices in the platform”. The custom-device constraint that survives applies to CL_DEVICE_TYPE_DEFAULT, which must not be a custom device unless it is the only device in the platform.
If you are carrying a decade-old mental model of that enum, it is out of date.
Profiles
Cutting across all of this is a distinction the article has so far skipped and most desktop developers never meet. Every platform and every device reports either FULL_PROFILE or EMBEDDED_PROFILE through CL_PLATFORM_PROFILE and CL_DEVICE_PROFILE.
The embedded profile is a formally specified relaxation of the full one, written for hardware that cannot afford the full set of guarantees: it permits round-to-zero instead of round-to-nearest as the default rounding mode for single precision, allows denormals and infinities to be handled more loosely, relaxes the accuracy requirements on several built-ins, drops 64-bit integers to optional, and lowers the minimum image dimensions and object counts an implementation must support.
This matters because “OpenCL 3.0 conformant” on a phone or an automotive SoC does not automatically mean the same arithmetic as OpenCL 3.0 conformant on a workstation.
It is also why Qualcomm’s statement that current Snapdragon platforms support OpenCL 3.0 full profile is a specific and meaningful claim rather than marketing throat-clearing. Check CL_DEVICE_PROFILE before you assume anything about numerics on embedded hardware.
Devices are grouped under platforms. A platform corresponds roughly to one vendor’s implementation. A single machine routinely exposes three or four: an Intel GPU platform, an NVIDIA platform, a PoCL CPU platform, a Rusticl platform. The mechanism that lets them coexist is the ICD, described in section 4.
Execution model
A kernel is executed over an NDRange: a 1-, 2-, or 3-dimensional index space. Each point in the index space is a work-item, and every work-item runs the same kernel body with a different global ID.
Work-items are grouped into work-groups of a size the application chooses (or leaves to the implementation).
Work-items in a work-group can synchronise with a barrier and share local memory; work-items in different work-groups cannot synchronise at all, and the specification gives no guarantee about the order in which work-groups execute or whether they execute concurrently.
That last point is the single most important constraint in the execution model and it is routinely violated by people porting from CPU code. There is no legal way to spin-wait in one work-group for another work-group to make progress. On many implementations it will appear to work and then deadlock on a device with fewer compute units or a different scheduler.
Since OpenCL 2.1 (and via cl_khr_subgroups before that) there is a third level: the sub-group, a subdivision of a work-group that maps onto the hardware’s SIMD execution width, a warp on NVIDIA, a wavefront on AMD, a subgroup on Intel and Arm and Qualcomm.
Sub-groups are guaranteed to make forward progress independently of one another within a work-group, and they support collective operations (broadcast, reduce, scan, ballot, shuffle) that are dramatically cheaper than the local-memory equivalents, and which became core in OpenCL 3.1.
Work is submitted through a command queue attached to one context and one device. Commands are kernel executions, memory transfers, map/unmap operations, markers, and barriers.
Each returns an event, and events form a dependency graph: any enqueue can take a wait-list of events that must complete first. Queues are in-order by default; CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE makes them out-of-order, at which point the event graph is the only thing constraining execution order.
The context is the object that ties devices, memory objects, programs, and queues together. Memory objects belong to a context, not a device, which is how the runtime knows it is allowed to migrate a buffer between devices in the same context.
Memory model
Four address spaces, in decreasing scope and increasing speed:
SpaceQualifierScopeTypical hardwareGlobal__globalAll work-items, all work-groups, host-visibleDevice DRAMConstant__constantRead-only, all work-itemsConstant cache / DRAMLocal__localOne work-groupScratchpad (LDS / shared memory)Private__privateOne work-itemRegisters, spilling to scratch
Since OpenCL 2.0 there is optionally a fifth: the generic address space, which lets a pointer be resolved to global, local, or private at runtime, so you can write one function that operates on all three.
Consistency is relaxed and explicit. Within a work-item, memory is consistent in program order. Across work-items in a work-group, memory is consistent only at a barrier(). Through work-groups, there is no consistency guarantee during a kernel’s execution at all, only at kernel boundaries and at explicit synchronisation points defined by the host API.
OpenCL 2.0 layered a formal memory model on top of this, adapted from C11, with atomic operations parameterised by memory order and memory scope (memory_scope_work_item, _work_group, _device, _all_svm_devices), which section 6 covers in detail.
Programming model
Two are named in the spec: data parallel (the NDRange) and task parallel (a kernel enqueued with a single work-item, or independent kernels in an out-of-order queue).
The task-parallel model is largely vestigial on GPUs; it mattered for the Cell processor and for CPU devices, and clEnqueueTask was deprecated in 2.0 in favour of just enqueuing a one-element NDRange.
The API, and what it feels like to use
OpenCL is a C API with an object model built on opaque handles and manual reference counting. Every object type has clRetain* and clRelease*.
Objects are freed when their count reaches zero, but the specification is careful to say that the implementation may keep an object alive as long as it is referenced by a queued command, so releasing a buffer that a running kernel still uses is legal.
The canonical bring-up sequence is longer than people expect:
#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
cl_platform_id platform;
cl_device_id device;
cl_int err;
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
/* clCreateCommandQueue is deprecated since 1.2; this is the 2.0+ form */
cl_queue_properties qprops[] = { CL_QUEUE_PROPERTIES,
CL_QUEUE_PROFILING_ENABLE, 0 };
cl_command_queue queue =
clCreateCommandQueueWithProperties(context, device, qprops, &err);
const char *src = "__kernel void saxpy(float a, \n"
" __global const float *x, \n"
" __global float *y) { \n"
" size_t i = get_global_id(0); \n"
" y[i] = fma(a, x[i], y[i]); \n"
"} \n";
cl_program program = clCreateProgramWithSource(context, 1, &src, NULL, &err);
err = clBuildProgram(program, 1, &device, "-cl-std=CL1.2", NULL, NULL);
if (err != CL_SUCCESS) {
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
0, NULL, &log_size);
char *log = malloc(log_size);
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
log_size, log, NULL);
fprintf(stderr, "%s\n", log);
}
cl_kernel kernel = clCreateKernel(program, "saxpy", &err);
cl_mem xbuf = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
n * sizeof(float), hx, &err);
cl_mem ybuf = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
n * sizeof(float), hy, &err);
float a = 2.0f;
clSetKernelArg(kernel, 0, sizeof(float), &a);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &xbuf);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &ybuf);
size_t global = n, local = 256;
cl_event ev;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, &ev);
clEnqueueReadBuffer(queue, ybuf, CL_TRUE, 0, n * sizeof(float), hy, 1, &ev, NULL);
The compiler is the core part of the runtime here: kernels are shipped as source strings and built at application startup, which is a design decision with large consequences.
And clSetKernelArg is stateful and untyped: the argument index and size are passed by hand, and getting either wrong is a runtime error at best and silent corruption at worst. This is the single biggest ergonomic difference from CUDA’s <<<>>> syntax, which type-checks arguments at compile time because the kernel and the launch site are in the same translation unit.
The C++ bindings (CL/opencl.hpp, formerly cl2.hpp, formerly cl.hpp) fix most of this with RAII wrappers and a variadic KernelFunctor. They are a Khronos-maintained header-only library, not part of the core spec, and they are what you should actually use from C++.
Error handling deserves a note. Every function returns or writes a cl_int, and there are around sixty error codes. Two behaviours cause most of the confusion:
Asynchronous errors don’t surface where you expect. A kernel that reads out of bounds does not produce an error from
clEnqueueNDRangeKernel; it producesCL_OUT_OF_RESOURCESor a device reset from a laterclFinish, or nothing at all. Attach a context error callback (thepfn_notifyparameter ofclCreateContext): most implementations report useful diagnostics through it and almost nobody sets it.CL_INVALID_WORK_GROUP_SIZEhas many causes. The local size must divide the global size (unless the device supports non-uniform work-groups, an OpenCL 2.0 feature that is optional in 3.0 and queried viaCL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT), must not exceedCL_KERNEL_WORK_GROUP_SIZEfor that specific kernel on that device, and must match anyreqd_work_group_sizeattribute in the kernel source.
Deprecated APIs remain callable. The headers hide 1.2-deprecated entry points unless you define CL_USE_DEPRECATED_OPENCL_1_2_APIS, and similarly for earlier versions.
The CL_TARGET_OPENCL_VERSION macro (a three-digit value: 120, 200, 300) controls which version’s prototypes the unified headers expose, and setting it deliberately is good hygiene: it turns “this device doesn’t support that” into the worst compile error.
The ICD, and how a call reaches a driver
Nothing in the OpenCL specification requires that multiple vendors’ implementations coexist on one machine, but in practice they must, and the mechanism is the Installable Client Driver loader.
An application links against libOpenCL.so (or OpenCL.dll, or libOpenCL.dylib). That library is not a driver. It is a dispatcher, maintained by Khronos in the OpenCL-ICD-Loader repository, which enumerates the vendor implementations installed on the system and forwards calls to the right one.
On Linux it does this by reading /etc/OpenCL/vendors/*.icd, each a one-line text file containing the path of a vendor’s shared object; on Windows it reads HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\OpenCL\Vendors.
Two environment variables control the search: OCL_ICD_VENDORS replaces the default directory, and OCL_ICD_FILENAMES takes a separator-delimited list of extra ICDs to load, which are enumerated before anything found by the default mechanism.
Both are invaluable when you want a test suite to hit one specific implementation on a machine that has four. Layers are selected separately, through OPENCL_LAYERS.
Mechanically, each vendor library exports clGetExtensionFunctionAddress and (per cl_khr_icd) clIcdGetPlatformIDsKHR. Every cl_platform_id the vendor returns points to a struct whose first member is a pointer to a dispatch table.
Every other OpenCL object (contexts, devices, queues, buffers) is likewise required to begin with a pointer to the dispatch table of the implementation that created it. So clEnqueueNDRangeKernel(queue, ...) in the loader is one indirect call through queue->dispatch->clEnqueueNDRangeKernel.
This is why passing objects from one platform into a call on another platform is undefined behaviour rather than a clean error: the loader has already dispatched before anyone checks.
Version 2.0.0 of cl_khr_icd added clIcdGetFunctionAddressForPlatformKHR and related entry points to handle extension functions cleanly, which the older scheme did badly.
The loader also supports layers, a mechanism borrowed from Vulkan. A layer is a shared object that sits between the application and the ICD and can intercept, log, modify, or synthesise OpenCL calls. This is the foundation of several genuinely useful tools:
OpenCL Intercept Layer (Intel, open source, version 3.0.4 as of 2026): call tracing, timing, kernel dumping, device-side printf capture, injection of modified kernel source without rebuilding the application, and USM validity checking. Runs on Windows, Linux, macOS, Android, and FreeBSD.
CLVizulayer (StreamHPC, presented at IWOCL 2026): emits the directed acyclic graph of device submissions in Graphviz DOT format. Unlike a timeline trace, this shows the constraints rather than the observed order, which is how you find the case where two commands ran sequentially because the implementation felt like it rather than because you asked for it. It has been used to inspect llama.cpp, Leela Chess Zero, GROMACS and LAMMPS.
Layers are underused. If you are debugging a portability problem across four implementations, the ability to insert instrumentation without touching either the application or the driver is worth a great deal.
The execution model in depth
NDRange geometry
clEnqueueNDRangeKernel takes work_dim (1–3), a global_work_offset (added in 1.1; usually NULL), global_work_size, and local_work_size.
Inside the kernel:
size_t gid = get_global_id(0); // global index, includes offset
size_t lid = get_local_id(0); // index within the work-group
size_t grp = get_group_id(0); // work-group index
size_t gsz = get_global_size(0);
size_t lsz = get_local_size(0); // actual, may differ from enqueued
size_t ngrp = get_num_groups(0);
size_t off = get_global_offset(0); // 1.1+
uint dim = get_work_dim();
Passing NULL for local_work_size lets the implementation choose. How much this matters is easy to measure, and the answer has two halves.
A 6.8× spread between the worst and best explicit choice. This is not a knob you can leave unconsidered; it is frequently the largest single factor in a kernel’s performance, ahead of most things people spend their time on.
And yet the implementation’s own choice beat every value I picked by hand. I’d written flatly that passing NULL is wrong for code you care about, and on this device that advice is (sorry) simply false: the runtime knows the vectorisation width its own compiler chose and I do not.
The honest version is much narrower: the parameter matters enormously, NULL is a real candidate rather than a lazy default, and the only way to know is to measure both on the hardware you ship to.
Where you do choose by hand, the size is bounded by three things you should query:
size_t max_wg, pref_mult;
cl_ulong local_used, private_used;
clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE,
sizeof(max_wg), &max_wg, NULL);
clGetKernelWorkGroupInfo(kernel, device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(pref_mult), &pref_mult, NULL);
clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(local_used), &local_used, NULL);
clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_PRIVATE_MEM_SIZE,
sizeof(private_used), &private_used, NULL);
CL_KERNEL_WORK_GROUP_SIZE is per-kernel, not per-device: it accounts for the register and local-memory pressure of this kernel and is therefore often much smaller than CL_DEVICE_MAX_WORK_GROUP_SIZE.
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE is the closest thing OpenCL 1.x has to a warp-size query: in practice it is 32 on NVIDIA, 64 on GCN and typically 32 on RDNA, 8, 16 or 32 on Intel depending on the SIMD width the compiler chose, and commonly 64 or 128 on Adreno.
Those are the values you will usually see, not values the specification promises; query it rather than hard-coding it, and a local size that is not a multiple of it wastes lanes.
CL_KERNEL_PRIVATE_MEM_SIZE is the register-spill indicator and the most useful single number for diagnosing a slow kernel. If it is non-zero and large, you are spilling to scratch memory and no amount of tuning the local size will help.
OpenCL 3.1 finally standardises what everyone had been approximating: a suggested local work-group size query, promoted to core from cl_khr_suggested_local_work_size. The runtime tells you what it thinks the right work-group size is for a given kernel, global size, and device. It is a hint, not an oracle, but it removes a class of per-device tuning tables from application code.
Sub-groups
uint sg_size = get_sub_group_size();
uint sg_id = get_sub_group_id();
uint sg_local = get_sub_group_local_id();
uint n_sg = get_num_sub_groups();
float total = sub_group_reduce_add(x);
float bcast = sub_group_broadcast(x, 0);
int any = sub_group_any(pred);
float shuf = sub_group_shuffle(x, src_lane); // cl_khr_subgroup_shuffle
float rot = sub_group_rotate(x, delta); // cl_khr_subgroup_rotate
uint4 mask = sub_group_ballot(pred); // cl_khr_subgroup_ballot
Sub-group operations avoid the round trip through local memory and the barrier that a work-group reduction requires. On a modern GPU a sub_group_reduce_add over 32 lanes is a handful of shuffle-and-add instructions; the local-memory equivalent is a store, a barrier, log₂(n) rounds of load-add-store with a barrier each, and a load.
For the reduction-heavy inner loops in attention and normalisation kernels this is the difference between competitive and not.
Before 3.1 this was a portability minefield: cl_khr_subgroups was optional, the extended type support was a separate extension (cl_khr_subgroup_extended_types), shuffles were another (cl_khr_subgroup_shuffle, cl_khr_subgroup_shuffle_relative), ballots another, and Intel had its own pre-standard cl_intel_subgroups.
The llama.cpp OpenCL backend requires subgroup support and therefore requires OpenCL 2.x or a 3.0 implementation that opted in. OpenCL 3.1 makes sub-groups core, including shuffles, rotations, and an expanded set of supported data types, which is the change most likely to simplify real kernel code.
You can pin the sub-group size with __attribute__((intel_reqd_sub_group_size(N))) on Intel, or query CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE and CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE via clGetKernelSubGroupInfo (2.1+).
Queues, events, and the flush trap
An event has five states: CL_QUEUED, CL_SUBMITTED, CL_RUNNING, CL_COMPLETE, or a negative value indicating abnormal termination. With CL_QUEUE_PROFILING_ENABLE set, clGetEventProfilingInfo returns nanosecond timestamps for CL_PROFILING_COMMAND_QUEUED, _SUBMIT, _START, _END, and (2.0+) _COMPLETE. END - START is device execution time; START - SUBMIT is queue latency; SUBMIT - QUEUED is host-side runtime overhead.
All three are worth looking at separately, a kernel that is fast but submitted late is a different problem from a kernel that is slow. The three intervals are easy to see. Timing a trivial kernel over 65,536 work-items with the queue warm:
Kernel work QUEUED→SUBMIT SUBMIT→START START → END
1 iteration 0.07 µs 1.75 µs 63.9 µs
50 iterations 0.13 µs 2.66 µs 3.6 ms
1000 iterations 0.37 µs 13.2 µs 89.0 ms
Host-side runtime cost is sub-microsecond, queue latency is single-digit microseconds, and everything else is the device. That ordering is what you want, and it is why the interesting overhead is not inside a single enqueue but in how many of them you make and how you make them.
Which is measurable too. Dispatching an empty kernel repeatedly, comparing a round trip per dispatch against submitting a batch and finishing once:
Dispatches Round-trip Batched Ratio
200 7.28 µs each 2.00 µs each 3.6×
1000 7.23 µs each 1.91 µs each 3.8×
5000 7.22 µs each 1.85 µs each 3.9×
PoCL 3.0, CPU device.
Nearly 4× on the same work, from nothing but submission discipline, and stable across three orders of magnitude of batch size. Note that this is measured against a CPU device with no PCIe bus and no kernel-mode driver round trip; on a discrete GPU the round-trip figure is worse, not better.
Around 2 µs of irreducible per-dispatch cost is also the number that makes command buffers interesting: a transformer decode step issuing three hundred kernels per token is paying roughly 600 µs of pure submission before any arithmetic happens, every token.
The classic OpenCL bug is forgetting clFlush. Enqueuing a command does not guarantee it is submitted to the device. If you enqueue work and then block on an event with clWaitForEvents without flushing the queue, some implementations will hang, because the command that would signal the event was never sent. The rules:
clFlushguarantees all previously enqueued commands are submitted to the device. It does not wait.clFinishblocks until all previously enqueued commands have completed. It implies a flush.Blocking enqueue calls on a queue (
clEnqueueReadBufferwithblocking_read = CL_TRUE) imply a flush of that queue.clWaitForEventsdoes not: the specification says the behaviour is undefined if you wait on events from commands that have not been flushed. This bites hardest across queues, since a blocking call on queue A flushes nothing in queue B. Flush explicitly.
OpenCL 3.1 fixes a genuinely subtle related hazard. Previously, polling an event’s status with clGetEventInfo and observing CL_COMPLETE did not by itself establish the memory ordering needed to safely read the results; you were expected to call a waiting function.
In practice a great deal of code polled and then read, and it mostly worked. In 3.1, observing that an event has reached CL_COMPLETE is itself a synchronisation point.
That is a spec change that legalises what people were already doing, which is the right call, but it means code written against 3.1 semantics can be subtly broken on a 3.0 driver.
Device-side enqueue, and why it did not take
OpenCL 2.0 introduced device-side enqueue: a kernel could enqueue further kernels onto a device-side queue without host involvement, using blocks (the Clang/Apple ^{} extension) as the payload.
It was the answer to CUDA Dynamic Parallelism, and it was intended for irregular workloads: adaptive mesh refinement, tree traversal, anything where the amount of work is data-dependent.
It never got broad adoption, unfortunately due to a lot of reasons. The implementation burden was high (it requires a device-side scheduler), the syntax was unfamiliar, and the performance on the implementations that did support it was frequently worse than doing multiple host-side dispatches.
In OpenCL 3.0 it became optional, queried through CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, and many implementations report zero.
It’s the clearest example of a 2.0 feature that was standardised before it was proven, which is exactly the mistake the working group says its current extension-first process is designed to avoid.
Command buffers
The replacement for a different problem, per-submission host overhead, is cl_khr_command_buffer, released provisionally in November 2021 as part of OpenCL 3.0.10 and developed largely by Codeplay with Qualcomm, Arm, Intel, Tampere University, NVIDIA and Google.
The idea is the one CUDA Graphs, Vulkan command buffers, and Level Zero command lists all landed on independently: record a sequence of commands once, finalise it, then dispatch the whole thing repeatedly with a single API call.
For inference serving, where the same graph of thirty or three hundred kernels is executed per token, the per-enqueue host cost dominates at small batch sizes.
cl_command_buffer_khr cb =
clCreateCommandBufferKHR(1, &queue, NULL, &err);
clCommandNDRangeKernelKHR(cb, NULL, NULL, k0, 1, NULL, &g0, &l0,
0, NULL, NULL, NULL);
clCommandNDRangeKernelKHR(cb, NULL, NULL, k1, 1, NULL, &g1, &l1,
0, NULL, NULL, NULL);
clFinalizeCommandBufferKHR(cb);
for (int i = 0; i < steps; i++)
clEnqueueCommandBufferKHR(0, NULL, cb, 0, NULL, NULL);
Note that a command buffer is internally out-of-order regardless of the queue it targets; ordering comes from the sync-point dependencies you declare when recording.
cl_khr_command_buffer_mutable_dispatch (OpenCL 3.0.12, September 2022) relaxes the immutability constraint so that kernel arguments, global size, local size and offsets of a recorded dispatch can be updated between replays via clUpdateMutableCommandsKHR, which is what you need for anything with a changing sequence length.
A further extension allows a command buffer to span multiple queues and devices (added in 3.0.14).
As of OpenCL 3.1 command buffers are still an extension, but Khronos names them explicitly as one of the features in the pipeline for a future core release.
The memory model in depth
Buffers, images, pipes
A buffer (cl_mem created by clCreateBuffer) is a linear byte range. An image (clCreateImage) is an opaque object with a channel order, channel data type, and dimensionality, accessed through read_imagef / write_imagef and friends, optionally through a sampler that performs normalised coordinates, addressing modes (clamp, repeat, mirror), and linear filtering in fixed-function hardware.
Pipes (2.0, optional in 3.0 via CL_DEVICE_PIPE_SUPPORT) are FIFO objects for producer-consumer patterns between kernels; they were designed with FPGAs in mind and are rarely used on GPUs.
The buffer-versus-image decision is more consequential than it looks. On desktop GPUs, buffer loads go through the general L1/L2 path; image reads go through the texture path, which on many architectures has a separate cache, hardware address computation, hardware boundary handling, and format conversion for free.
On mobile GPUs, Adreno in particular, the texture path is substantially faster for the read patterns typical of convolution and GEMM. This is why Qualcomm’s TVM and MLC work for Adreno has a dedicated “texture path” and specialised layouts, and why the llama.cpp Adreno kernels are written the way they are. On CPU devices the distinction mostly evaporates and images are usually slower.
Image support is itself optional: CL_DEVICE_IMAGE_SUPPORT can be CL_FALSE. It commonly is on custom devices and on some early open-source stacks, the original Mesa Clover never supported images, which is precisely why it could not run darktable, and why Rusticl’s image support was the thing that made it useful.
Image formats and samplers
An image is not a typed buffer; it is a channel order plus a channel data type, and only some combinations are required. The channel orders in the specification are CL_R, CL_A, CL_RG, CL_RA, CL_RGB, CL_RGBA, CL_BGRA, CL_ARGB, CL_INTENSITY, CL_LUMINANCE, CL_DEPTH and CL_sRGBA, among others.
The data types run from CL_SNORM_INT8 and CL_UNORM_INT8 through the packed short formats (CL_UNORM_SHORT_565, CL_UNORM_SHORT_555, CL_UNORM_INT_101010) to CL_SIGNED_INT8/16/32, CL_UNSIGNED_INT8/16/32, CL_HALF_FLOAT and CL_FLOAT, with CL_UNORM_INT10, INT12 and INT14 added for higher-precision sensor data.
The normalised types do conversion in hardware: reading a CL_UNORM_INT8 image with read_imagef returns floats in [0, 1] with no instruction spent on the divide. That is free range conversion, and a real reason to prefer images for pixel data.
Only a small subset of order and type combinations is guaranteed, though; everything else must be checked with clGetSupportedImageFormats against the specific device, memory flags and image type. Assuming a format exists because it is in the enum is one of the more common portability failures.
Samplers carry three orthogonal settings: normalised or unnormalised coordinates, an addressing mode for out-of-range coordinates (CLK_ADDRESS_NONE, CLAMP, CLAMP_TO_EDGE, REPEAT, MIRRORED_REPEAT), and a filter mode (CLK_FILTER_NEAREST or CLK_FILTER_LINEAR). All of it is fixed-function on GPUs.
A bilinear tap that would cost four loads and three lerps in a buffer kernel is one read_imagef with CLK_FILTER_LINEAR, and boundary clamping that would cost a branch per axis is free.
Samplers can be declared in the kernel as a const sampler_t constant or created on the host and passed in.
Allocation flags and the map path
CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY /* access, from the kernel's view */
CL_MEM_USE_HOST_PTR /* use this host allocation */
CL_MEM_ALLOC_HOST_PTR /* allocate host-accessible memory */
CL_MEM_COPY_HOST_PTR /* allocate device memory, copy in */
CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS /* 1.2 */
The semantics people get wrong: CL_MEM_USE_HOST_PTR doesn’t mean “the device will read your pointer directly”.
It means the implementation may cache the contents in device memory and is required to keep the host pointer as the backing store; you must map/unmap to access it safely from the host.
CL_MEM_ALLOC_HOST_PTR is the closest OpenCL comes to CUDA’s pinned memory: it asks the implementation to allocate memory the host can access efficiently, which on a discrete GPU usually means page-locked system memory suitable for DMA, and on an integrated GPU usually means memory both processors can access without a copy at all.
It is worth being precise about what the flag buys, because it is easy to assume it makes transfers faster on its own. Writing 64 MB into a buffer three ways on the same device:
Path Time Eff. Rate clEnqueueWriteBuffer into a plain buffer 14.9 ms 4.5 GB/s
clEnqueueWriteBuffer into an ALLOC_HOST_PTR buffer 29.1 ms 2.3 GB
/smap / write / unmap on the ALLOC_HOST_PTR buffer 5.2 ms 12.9 GB/s
PoCL 3.0, CPU device, mean of five.
Adding CL_MEM_ALLOC_HOST_PTR and changing nothing else made the copy twice as slow. The 2.9× win only appears when the access pattern changes to match the allocation: when you stop copying and start writing directly into mapped memory.
The flag is not an optimisation; it is a request for memory with different properties, and it pays only if you then use those properties. This generalises: allocation hints in OpenCL are contracts about placement, and a contract you do not exercise is overhead.
The idiomatic zero-copy pattern on integrated hardware is therefore CL_MEM_ALLOC_HOST_PTR plus clEnqueueMapBuffer:
cl_mem b = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
n, NULL, &err);
float *p = clEnqueueMapBuffer(queue, b, CL_TRUE, CL_MAP_WRITE,
0, n, 0, NULL, NULL, &err);
/* fill p */
clEnqueueUnmapMemObject(queue, b, p, 0, NULL, NULL);
On a system where host and device share physical memory this can be a genuine no-copy path. On a discrete GPU it still copies, but from pinned memory, at full PCIe bandwidth rather than the roughly half you get from pageable memory.
OpenCL 3.1 clarifies the semantics of CL_DEVICE_HOST_UNIFIED_MEMORY so that it can be used reliably to distinguish integrated from discrete devices.
That sounds trivial, but it’s not: for a decade the flag was defined loosely enough that implementations disagreed about it, and every serious application ended up with its own heuristic (usually string-matching the device name) to decide whether to take the map path or the copy path, so having one query that means one thing removes real code from real applications.
clEnqueueMigrateMemObjects (1.2) lets you explicitly move a memory object to a device, or to the host with CL_MIGRATE_MEM_OBJECT_HOST, ahead of the kernel that will use it. In multi-device contexts this is how you avoid a fault-driven migration in the middle of a dispatch.
The 2.0 memory model
OpenCL 1.x had no formal memory model, just “relaxed consistency“ plus barriers and fences, described in prose. OpenCL 2.0 replaced this with a model derived from C11’s, which was the right decision and made OpenCL one of the first GPU APIs with a mathematically specified memory model.
Atomics are typed (atomic_int, atomic_float, atomic_uintptr_t, and atomic_long/atomic_double if supported) and take a memory order and a memory scope:
atomic_fetch_add_explicit(&counter, 1,
memory_order_relaxed,
memory_scope_work_group);
atomic_store_explicit(&flag, 1,
memory_order_release,
memory_scope_device);
Orders are relaxed, acquire, release, acq_rel, seq_cst. Scopes are work_item, sub_group, work_group, device, and all_svm_devices. The scope is the part with no C11 analogue and it is what makes the model usable on GPUs: an atomic scoped to a work-group can be implemented in the scratchpad with no cache-coherence traffic, while an atomic scoped to all_svm_devices may require flushing through to system memory.
The historical wrinkle is the inclusive scope rule. Until 2026, for two atomic operations to synchronise with each other, their scopes had to match, or more precisely, both had to include each other’s work-items.
This essentially meant that a release at memory_scope_work_group did not necessarily synchronise with an acquire at memory_scope_device, even though the device scope is strictly larger.
It is a rule that made implementations easier and reasoning harder, and it produced code that specified everything at device scope out of caution, giving up the performance that scoped atomics exist to provide.
OpenCL 3.1 relaxes it: scopes no longer have to match exactly, and a finer-grained scope can satisfy a coarser-grained synchronisation requirement.
Combined with sub-groups becoming core, this makes the fine-grained synchronisation patterns used in modern GPU kernels expressible without the previous defensive over-specification.
Shared Virtual Memory, and the USM successor
OpenCL 2.0 added SVM in three tiers, queried through CL_DEVICE_SVM_CAPABILITIES:
Coarse-grained buffer: CL_DEVICE_SVM_COARSE_GRAIN_BUFFER. Pointers are valid on both sides; consistency at map/unmap and kernel boundaries.
Fine-grained buffer: CL_DEVICE_SVM_FINE_GRAIN_BUFFER. Concurrent host/device access to the same allocation, no map needed.
Fine-grained system: CL_DEVICE_SVM_FINE_GRAIN_SYSTEM. Any malloc pointer is usable on the device.
Atomics: CL_DEVICE_SVM_ATOMICS. Cross-side atomics on SVM allocations.
Coarse-grained buffer SVM is widely supported. Fine-grained system SVM is rare and requires real hardware support for demand paging and coherent access to host page tables.
The problem with SVM is not the concept but the API. clSVMAlloc returns a pointer, but you still set it as a kernel argument through clSetKernelArgSVMPointer, you still have to declare indirect accesses with clSetKernelExecInfo and CL_KERNEL_EXEC_INFO_SVM_PTRS, and the capability matrix is coarse.
Intel’s answer, cl_intel_unified_shared_memory (added to the registry alongside OpenCL 3.0.10), is a cleaner design borrowed from what became SYCL 2020’s USM: three allocation kinds (host, device, and shared) with explicit control over placement and migration, no map/unmap, allocations associated with both a device and a context, and a richer capability query.
It is the model that oneAPI, SYCL, and by extension a lot of HPC code actually use.
The working group has been standardising this as cl_khr_unified_svm, and Khronos lists Unified Shared Memory as one of the extensions in flight for a future core version.
For anyone writing new OpenCL against Intel or Level Zero-adjacent stacks, USM is already the pragmatic choice; for portable code, SVM coarse-grained remains the lowest common denominator, and plain buffers remain the actual lowest common denominator.
OpenCL C
OpenCL C is C99 with removals and additions. Removed: function pointers, recursion, variable-length arrays, bit fields, most of the standard library, goto into blocks, and (before 2.0) program-scope variables in the global address space.
Added: address space qualifiers, vector types, a large built-in function library, work-item query functions, and a set of type qualifiers and attributes.
Vector types exist in widths 2, 3, 4, 8, and 16 for all scalar types: float4, int8, uchar16, double2. They support arithmetic elementwise, and component access through several syntaxes:
float4 v = (float4)(1.0f, 2.0f, 3.0f, 4.0f);
float a = v.x; /* also .y .z .w; and .r .g .b .a since OpenCL 3.0 */
float2 b = v.xy;
float4 c = v.wzyx; /* arbitrary swizzle, including repeats */
float d = v.s3; /* hex-index form, required for width 8 and 16 */
float8 e = (float8)(v, v);
float2 lo = v.lo, hi = v.hi; /* halves */
float2 ev = v.even, od = v.odd;
Three-component vectors have size 16 bytes, not 12: sizeof(float3) == sizeof(float4). This surprises people writing structs that cross the host/device boundary.
Whether vector types help performance is architecture-dependent, and the answer has inverted over time. On AMD’s pre-GCN VLIW architectures (TeraScale) and on CPU devices, vector code was essential because the compiler could not always find the parallelism itself.
On GCN, RDNA, NVIDIA, and modern Intel GPUs, the hardware is scalar-per-lane and the vector types are mostly a way to express wide loads and stores. That is still worth something: a float4 load is one 128-bit memory instruction instead of four 32-bit ones, which matters for memory-bound kernels.
On Adreno and Mali, vector width still maps to real SIMD capability and choosing it correctly matters more.
Precision, and the parts nobody reads
The specification contains a table of maximum error in ULP for every math built-in, and it is the most under-appreciated part of the document. Some entries:
Function Max error
x + y, x * y, fma correctly rounded
1.0/x, x / y ≤ 2.5 ULP; correctly rounded with -cl-fp32-correctly-rounded-divide-sqrt
sqrt ≤ 3 ULP; correctly rounded with the same flag
rsqrt, cbrt, log1p ≤ 2 ULP
exp, exp2, exp10,
log, log2, log10 ≤ 3 ULP
sin, cos, sinpi,
cospi, hypot ≤ 4 ULP
tan, tanh, atan, atanh ≤ 5 ULP
pow, pown, powr, rootn,
erf, erfc, tgamma ≤ 16 ULP
mad unbounded: any value is conforming
native_* variants implementation-defined, no bound
half_* variants ≤ 8192 ULP
Two entries deserve attention. mad has no accuracy requirement at all: the specification permits any result, because it exists to let the implementation pick whatever multiply-add the hardware has, fused or not.
If you want a fused multiply-add with defined semantics, write fma. And in double precision, division, reciprocal and square root are all correctly rounded, the loose bounds above are a single-precision phenomenon.
What the table does not tell you is what you will actually get, and the gap between the two is larger than most people assume. Measuring the observed error of a conformant implementation against a double-precision reference over 65,536 points per function.
Two results are worth sitting with. native_sqrt, native_exp and native_log returned bit-identical values to their accurate counterparts: zero of 65,536 samples differed in any bit, so this implementation simply aliases them.
native_sin is the exception, differing on 21% of samples. half_exp, permitted to be wrong by 8192 ULP, was accurate to under one. So on this implementation, every fast-math variant is free accuracy, and code written to tolerate 8192 ULP is running on results good to 1.
That is not a reassuring finding. It means you cannot learn anything about native_ behaviour by testing it, because the next implementation is equally entitled to return something wildly different and still be conformant.
This is the article’s opening constraint in miniature: a standard spanning unknown hardware can specify floors and nothing else, and a floor tells you almost nothing about the room. Test on the implementation you ship against, or use the accurate functions.
The native_ family (native_sin, native_exp2, native_recip, native_rsqrt, native_divide) maps directly to whatever hardware instruction exists and gives no accuracy guarantee at all.
On most GPUs these are single-instruction and roughly an order of magnitude faster than the accurate versions. -cl-fast-relaxed-math implies -cl-finite-math-only, -cl-unsafe-math-optimizations, and permits the compiler to substitute native_ variants globally, which is why enabling it can change results by several ULP and can turn a NaN check into dead code.
Denormal handling is implementation-defined for single precision unless the device reports CL_FP_DENORM in CL_DEVICE_SINGLE_FP_CONFIG. -cl-denorms-are-zero explicitly permits flush-to-zero.
Most GPUs flush single-precision denormals by default and handle double-precision denormals correctly, which is the opposite of what people assume.
Double precision requires cl_khr_fp64 (or __opencl_c_fp64 in OpenCL C 3.0), half precision requires cl_khr_fp16. Neither is guaranteed. half as a storage type, via vload_half / vstore_half, which convert to and from float, is available without cl_khr_fp16; only arithmetic on half requires the extension.
This is a useful distinction for anyone storing FP16 weights and computing in FP32.
Attributes and hints
__attribute__((reqd_work_group_size(16, 16, 1)))
__attribute__((work_group_size_hint(64, 1, 1)))
__attribute__((vec_type_hint(float4)))
__attribute__((intel_reqd_sub_group_size(16))) /* vendor */
reqd_work_group_size is a contract: enqueue with any other local size and you get CL_INVALID_WORK_GROUP_SIZE.
It is worth using, because it lets the compiler size local arrays statically, unroll fully, and allocate registers knowing the occupancy, which frequently produces measurably better code than the hint version.
OpenCL C 3.0 and feature macros
OpenCL 3.0 made almost everything from the 2.x line optional. In the language, that optionality is expressed through predefined feature-test macros named __opencl_c_<feature>, which the compiler defines with value 1 when the feature is present:
__opencl_c_3d_image_writes
__opencl_c_atomic_order_acq_rel
__opencl_c_atomic_order_seq_cst
__opencl_c_atomic_scope_device
__opencl_c_atomic_scope_all_devices
__opencl_c_device_enqueue
__opencl_c_fp64
__opencl_c_generic_address_space
__opencl_c_images
__opencl_c_int64
__opencl_c_pipes
__opencl_c_program_scope_global_variables
__opencl_c_read_write_images
__opencl_c_subgroups
__opencl_c_work_group_collective_functions
So portable OpenCL C 3.0 looks like this:
#if defined(__opencl_c_subgroups)
float total = sub_group_reduce_add(partial);
#elif defined(__opencl_c_work_group_collective_functions)
float total = work_group_reduce_add(partial);
#else
/* hand-rolled local memory tree reduction */
#endif
On the host side there is a matching set of device queries: CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, CL_DEVICE_ATOMIC_FENCE_CAPABILITIES, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, CL_DEVICE_PIPE_SUPPORT, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT, CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT, CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, and CL_DEVICE_OPENCL_C_ALL_VERSIONS / CL_DEVICE_OPENCL_C_FEATURES, so an application can decide which kernel variant to build before it builds anything. clang exposes the same switches for offline compilation through -cl-ext, e.g. -cl-std=CL3.0 -cl-ext=+cl_khr_fp64,+__opencl_c_fp64.
C++, twice
OpenCL 2.2 (May 2017) introduced OpenCL C++, a static subset of C++14 as a kernel language, defined alongside SPIR-V ingestion. It is a carefully written specification. It appears never to have shipped in a production driver, there is no extension defined to detect support for it, and it was deprecated in OpenCL 3.0.
The replacement is C++ for OpenCL, which is a different thing: not a Khronos-ratified specification but a community language, documented in the OpenCL-Docs repository and implemented in upstream Clang since release 9.
It is C++17 layered over OpenCL C, backward compatible with OpenCL C source, and it works with any implementation that ingests SPIR-V, which, as of OpenCL 3.1, is all of them. Version 1.0 was published in December 2020 (compatible with OpenCL 2.0); the 2021 revision (December 2021) is compatible with OpenCL 3.0.
What it doesn’t support: virtual functions, dynamic_cast, non-placement new/delete, exceptions, pointers to member functions, references to functions, and the C++ standard library.
What it does support is templates, classes, operator overloading, lambdas, and auto, which is enough for the thing people actually want, namely writing one templated kernel body instead of four macro-expanded copies.
Address spaces are extended across C++ constructs: functional casts, templates, class members, references, lambdas, and operators.
template<typename T>
class complex_t {
T re, im;
public:
complex_t(T r, T i) : re{r}, im{i} {}
complex_t operator*(const complex_t &o) const {
return { re*o.re - im*o.im, re*o.im + im*o.re };
}
T real() const { return re; }
T imag() const { return im; }
};
__kernel void mul_sp(__global float *in, __global float *out) {
auto i = get_global_id(0);
auto r = complex_t{in[4*i], in[4*i+1]} * complex_t{in[4*i+2], in[4*i+3]};
out[2*i] = r.real(); out[2*i+1] = r.imag();
}
Online compilation of C++ for OpenCL requires the cl_ext_cxx_for_opencl extension and -cl-std=CLC++; Arm announced support in December 2020. In practice most people compile it offline with Clang to SPIR-V, which is the flow the language was designed for.
Compilation: source, IR, and machine code
The four ways to make a program object
clCreateProgramWithSource(ctx, count, strings, lengths, &err);
clCreateProgramWithIL(ctx, il, length, &err); /* 2.1+, SPIR-V */
clCreateProgramWithBinary(ctx, n, devs, sizes, bins, st, &err);
clCreateProgramWithBuiltInKernels(ctx, n, devs, names, &err); /* 1.2+ */
Source is the original path and still the most common. It means the OpenCL driver contains a full C compiler, that compiler runs at application startup, and compile time is user-visible. A large kernel library (llama.cpp’s OpenCL backend, a serious image processing pipeline) can take seconds to build.
It is worth knowing the size of that cost rather than assuming it. Building a synthetic kernel library on PoCL’s CPU device, then rebuilding the same library from the binaries the first build produced:
Kernels in the librarySource sizeclCreateProgramWithSource + buildclCreateProgramWithBinary + buildRatio83.5 KB158.0 ms2.2 ms72×3214.1 KB190.4 ms5.2 ms37×6428.2 KB276.4 ms13.8 ms20×
The absolute numbers are one implementation’s, and PoCL is running a full Clang and LLVM pipeline where a vendor driver would be more streamlined.
The shape is the point: compiling from source is 20× to 72× more expensive than loading a binary here, and the gap narrows only slowly as the library grows because a large part of the source cost is fixed compiler startup.
For an application with a hundred kernels this is the difference between an instant launch and a visible pause, which is exactly why every serious OpenCL application ends up building a cache, and exactly what the OpenCL 3.1 SPIR-V mandate is for.
Binaries via clCreateProgramWithBinary are the obvious fix and come with a hard constraint: program binaries are not portable. They are opaque, implementation-defined blobs, valid only for the device, driver version, and build options that produced them. CL_PROGRAM_BINARY_TYPE tells you whether a binary is an executable, a compiled object, or a library.
The correct use is a persistent cache keyed on device name, driver version, platform version, kernel source hash, and build options, exactly what llama.cpp does with GGML_OPENCL_KERNEL_CACHE_DIR, defaulting to %LOCALAPPDATA%\llama.cpp\cl-cache on Windows and ~/Library/Caches/llama.cpp/cl-cache on macOS, and invalidating on any of those inputs changing.
Separate compilation and linking arrived in 1.2:
clCompileProgram(prog, n, devs, opts, n_hdrs, hdrs, hdr_names, cb, data);
cl_program lib = clLinkProgram(ctx, n, devs, "-create-library",
n_in, inputs, cb, data, &err);
This lets you build a device-side library once and link it into several programs, and it lets you pass headers by name rather than concatenating strings.
Build options worth knowing:
-cl-std=CL1.2 | CL2.0 | CL3.0 | CLC++ | CLC++2021
-D name=value -I dir
-cl-single-precision-constant -cl-denorms-are-zero
-cl-opt-disable -cl-mad-enable
-cl-no-signed-zeros -cl-unsafe-math-optimizations
-cl-finite-math-only -cl-fast-relaxed-math
-cl-fp32-correctly-rounded-divide-sqrt
-cl-uniform-work-group-size -cl-kernel-arg-info
-w -Werror
-cl-kernel-arg-info is the one people forget: without it, clGetKernelArgInfo cannot tell you argument names, types, or address space qualifiers, which breaks any tooling that wants to bind arguments by name.
SPIR, SPIR-V, and the 2026 mandate
The first attempt at a portable IR was SPIR (2012, versions 1.2 and 2.0), which was LLVM IR with an OpenCL-specific metadata layer. It inherited LLVM’s problem: LLVM IR is not a stable format and its semantics are defined by the LLVM version that produced it. Consuming SPIR meant, in practice, having a compatible LLVM inside the driver.
SPIR-V, released with OpenCL 2.1 in November 2015 and adopted by Vulkan 1.0 three months later, replaced it with a purpose-built, versioned, SSA-form binary IR that is not tied to any compiler. It is shared with Vulkan, and a SPIR-V module declares which execution environment it targets, the OpenCL SPIR-V Environment Specification defines what an OpenCL implementation must accept.
The OpenCL flavour uses the OpenCL memory model and the Kernel execution model, and calls into the OpenCL Extended Instruction Set for SPIR-V for math built-ins.
There are now several ways to produce it:
Clang’s SPIR-V target.
clang -target spirv64 -c kernel.cl -o kernel.spv, using the SPIR-V backend that has been maturing in upstream LLVM. This is the path that no longer requires a translator.SPIRV-LLVM-Translator, the older, still widely used tool that converts LLVM IR to SPIR-V; the standard route for
clang -cl-std=CL3.0 -emit-llvmoutput and for oneAPI’s toolchain.clspv. Google’s compiler from a subset of OpenCL C to Vulkan compute shaders, which is a different SPIR-V dialect. Paired with clvk, a runtime that implements the OpenCL API on top of Vulkan.
OpenCL 3.1’s headline change is that SPIR-V ingestion is mandatory. Every conformant OpenCL 3.1 implementation must accept SPIR-V kernels through clCreateProgramWithIL, and must additionally support the SPIR-V query extension so applications can enumerate which SPIR-V capabilities, extensions, and versions a device handles.
This is more consequential than it sounds, and it is worth being precise about why. Before 3.1, clCreateProgramWithIL was core in 2.1 but optional in 3.0, so a 3.0 implementation could legally accept only source. Any tool that wanted to target OpenCL as a backend therefore had to either ship an OpenCL C source generator or accept that it would not run everywhere.
That is a real tax on SYCL implementations, on chipStar (which compiles CUDA and HIP to SPIR-V), on Julia’s and Rust’s GPU backends, and on every domain-specific compiler. Making ingestion mandatory turns OpenCL from “an API you can target if the driver cooperates” into a guaranteed compilation target.
Neil Trevett, who chairs the working group, called it the most consequential change in 3.1, and on the evidence that is not marketing.
The secondary benefits are the ones that matter operationally: kernels can ship pre-compiled and pre-optimised rather than as source, which removes startup compile cost, allows ahead-of-time specialisation, and means you are not shipping your kernel source to customers.
What the vendor stacks actually do
ImplementationFront endIRBack endNVIDIAClang-derivedNVVM (LLVM IR)PTX → SASS via ptxasIntel (NEO / compute-runtime)ClangSPIR-VIGC → Gen ISAAMD (ROCm CLR)ClangLLVM IRAMDGPU back end → GCN/RDNA ISAArm MaliClang-basedvendor IRMali ISAQualcomm AdrenoLLVM-basedvendor IRAdreno ISAPoCLClangLLVM IR / SPIR-VLLVM target back ends, plus CUDA / Level Zero / remote driversRusticlvia SPIRV-Tools / MesaSPIR-V → NIRGallium driver back endsclvk / clspvclspv (Clang-based)Vulkan-flavour SPIR-Vwhatever the Vulkan driver does
Apple’s implementation, NVIDIA’s, AMD’s, RapidMind’s and Gallium’s have all been LLVM-based from early on.
The practical implication is that “OpenCL C” in the field means “whatever Clang’s OpenCL front end accepts, plus vendor quirks” more than it means the paper specification, which is mostly good, because Clang’s OpenCL support is well maintained and tracks the spec closely, and occasionally bad, because a kernel that builds on four Clang-derived stacks may still fail on a non-LLVM one.
Performance engineering
The mental model
The abstraction says work-items are independent. The hardware executes them in lockstep groups. Everything about OpenCL performance follows from taking the second statement seriously while writing code that satisfies the first.
Concretely: a work-group is scheduled onto one compute unit and stays there. It is subdivided into sub-groups, which are the actual scheduling and execution unit. Divergent control flow within a sub-group is executed by predicating both sides.
Memory requests from a sub-group are coalesced by the hardware into as few transactions as possible; consecutive lanes reading consecutive addresses is one transaction, consecutive lanes reading strided addresses is many.
The occupancy story is the same as CUDA’s without the vocabulary: a compute unit has a fixed register file and a fixed scratchpad, and the number of work-groups it can host concurrently is bounded by whichever runs out first.
More concurrent work-groups means more latency hiding. CL_KERNEL_LOCAL_MEM_SIZE and CL_KERNEL_PRIVATE_MEM_SIZE are how you see the pressure; CL_DEVICE_LOCAL_MEM_SIZE and CL_DEVICE_MAX_WORK_GROUP_SIZE are the budget.
A worked GEMM
The single-work-item-per-output-element version is the reference, and it is memory bound at an arithmetic intensity of about 1 FLOP per byte:
__kernel void sgemm_naive(const int M, const int N, const int K,
__global const float *A,
__global const float *B,
__global float *C)
{
const int col = get_global_id(0);
const int row = get_global_id(1);
float acc = 0.0f;
for (int k = 0; k < K; k++)
acc += A[row*K + k] * B[k*N + col];
C[row*N + col] = acc;
}
Tiling through local memory raises intensity by the tile width. Each work-group cooperatively stages a TS × TS tile of A and of B, barriers, then every work-item does TS multiply-accumulates out of the scratchpad:
#define TS 16
__attribute__((reqd_work_group_size(TS, TS, 1)))
__kernel void sgemm_tiled(const int M, const int N, const int K,
__global const float *A,
__global const float *B,
__global float *C)
{
const int lx = get_local_id(0), ly = get_local_id(1);
const int col = get_group_id(0)*TS + lx;
const int row = get_group_id(1)*TS + ly;
__local float Asub[TS][TS];
__local float Bsub[TS][TS];
float acc = 0.0f;
for (int t = 0; t < K/TS; t++) {
Asub[ly][lx] = A[row*K + (t*TS + lx)];
Bsub[ly][lx] = B[(t*TS + ly)*N + col];
barrier(CLK_LOCAL_MEM_FENCE);
#pragma unroll
for (int k = 0; k < TS; k++)
acc = fma(Asub[ly][k], Bsub[k][lx], acc);
barrier(CLK_LOCAL_MEM_FENCE);
}
C[row*N + col] = acc;
}
TS is 16 here rather than the textbook 32 for a reason: a 32×32 work-group is 1024 work-items, which exceeds CL_DEVICE_MAX_WORK_GROUP_SIZE on a great deal of mobile and embedded hardware and sits at the ceiling on most desktop GPUs. Both kernels here also assume M, N and K are multiples of TS; real code needs edge handling or padded allocations.
Both barriers are required. The second one, after the inner loop and before the next tile overwrites the scratchpad, is the one people omit, and it produces a race that is invisible on any implementation where a work-group happens to be one sub-group wide.
The tiled version is usually still short of peak, because each work-item does one FMA per two scratchpad reads. The next step is register tiling: give each work-item a WPT × WPT block of outputs so operands loaded into registers are reused across several accumulations.
#define TS 32 /* tile size */
#define WPT 4 /* outputs per work-item per dimension */
#define RTS (TS/WPT)
__attribute__((reqd_work_group_size(RTS, RTS, 1)))
__kernel void sgemm_regtiled(const int M, const int N, const int K,
__global const float *A,
__global const float *B,
__global float *C)
{
const int lx = get_local_id(0), ly = get_local_id(1);
const int gx = get_group_id(0)*TS, gy = get_group_id(1)*TS;
__local float Asub[TS][TS];
__local float Bsub[TS][TS];
float acc[WPT][WPT];
#pragma unroll
for (int a = 0; a < WPT; a++)
#pragma unroll
for (int b = 0; b < WPT; b++) acc[a][b] = 0.0f;
for (int t = 0; t < K/TS; t++) {
#pragma unroll
for (int a = 0; a < WPT; a++)
#pragma unroll
for (int b = 0; b < WPT; b++) {
const int r = ly*WPT + a, c = lx*WPT + b;
Asub[r][c] = A[(gy + r)*K + t*TS + c];
Bsub[r][c] = B[(t*TS + r)*N + gx + c];
}
barrier(CLK_LOCAL_MEM_FENCE);
#pragma unroll
for (int k = 0; k < TS; k++) {
float av[WPT], bv[WPT];
#pragma unroll
for (int a = 0; a < WPT; a++) av[a] = Asub[ly*WPT + a][k];
#pragma unroll
for (int b = 0; b < WPT; b++) bv[b] = Bsub[k][lx*WPT + b];
#pragma unroll
for (int a = 0; a < WPT; a++)
#pragma unroll
for (int b = 0; b < WPT; b++)
acc[a][b] = fma(av[a], bv[b], acc[a][b]);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
#pragma unroll
for (int a = 0; a < WPT; a++)
#pragma unroll
for (int b = 0; b < WPT; b++)
C[(gy + ly*WPT + a)*N + gx + lx*WPT + b] = acc[a][b];
}
Now WPT loads of A and WPT of B feed WPT² FMAs. With WPT = 4 that is 16 FMAs per 8 scratchpad reads.
All three kernels above were compiled and run before publication, on PoCL 3.0 targeting a Xeon CPU device, and checked against a NumPy reference at 256×256: all three match to zero relative error. That is a correctness check, not a performance claim, and the CPU result below is the reason to keep those two things apart.
There is no single correct TS and WPT. The optimum depends on scratchpad size, register file size, sub-group width, cache line size, and whether the compiler decides to spill. This is exactly what CLBlast does, a tuned BLAS in OpenCL that ships per-architecture parameter sets and can autotune on unseen hardware, and what CLTune and KTT exist to automate.
The honest position is that hand-written OpenCL GEMM is a tuning problem, not a coding problem, and that any claim about “OpenCL performance” that does not say which parameters were used is not a measurement.
Scratchpad bank conflicts are the other thing to watch. Local memory is divided into banks (typically 32 words wide); simultaneous accesses by different lanes to different addresses in the same bank serialise.
Declaring __local float Asub[TS][TS+1], padding the row stride by one, breaks the conflicting stride pattern at a cost of a few hundred bytes.
Everything that isn’t the kernel
For inference workloads the kernel is frequently not the bottleneck. Intel’s IWOCL 2026 talk on optimising AI workloads on their GPUs spent most of its time on three things that are not kernel code: queue model design and how to batch versus submit immediately, Ultra Low Latency Submission to cut dispatch overhead for latency-sensitive inference, and driver-level memory pooling and resource recycling to avoid the allocation churn that dominates when a model allocates and frees repeatedly.
They reported allocation overhead reductions of orders of magnitude from pooling alone.
That matches what anyone who has profiled a token-by-token decode loop finds: at batch size 1, per-dispatch host cost and allocator behaviour can be a larger share of wall time than the arithmetic. It is why command buffers exist, and why they are the extension most worth watching.
Asynchronous copies
For embedded and DSP targets where the scratchpad is filled by an explicit DMA engine rather than by ordinary loads, OpenCL C has:
event_t e = async_work_group_copy(dst_local, src_global, n, 0);
event_t f = async_work_group_strided_copy(dst_local, src_global, n, stride, e);
wait_group_events(1, &f);
prefetch(ptr, n);
cl_khr_extended_async_copies and cl_khr_async_copy_fence, added in OpenCL 3.0.10, extend this with 2D and 3D copy patterns and finer-grained fencing.
These were introduced specifically for the class of embedded processors that motivated much of the 3.0 work, and on a GPU they are usually implemented as ordinary loads.
Performance portability, honestly
The two studies everyone quotes are worth reading rather than citing, because the numbers travel without their conditions.
Karimi, Dickson and Hamze at D-Wave measured a quantum Monte Carlo kernel in both languages and reported the OpenCL kernel between about 13% and 63% slower, with end-to-end times 16% to 67% slower. Those figures are exact, and they are also from 2010, on a GeForce GTX-260, with the CUDA and OpenCL toolkits both at version 2.3.
More importantly, the paper is explicit that its OpenCL kernel is a near-identical port of the CUDA one, changed only where OpenCL forced a change: __shared__ to __local, threadIdx to get_local_id(), __syncthreads() to barrier(), and one array-indexing rewrite.
It measures what a direct translation costs, which is a real and useful thing to know, and it doesn’t measure what tuned OpenCL costs. The spread also narrows as problems get larger, which the authors attribute to the kernel’s share of total runtime rising.
The 2011 Delft comparison is usually summarised as CUDA leading a straightforward OpenCL translation by at most 30% on NVIDIA hardware, with the gap attributed to programming-model differences and to NVIDIA having invested more in its CUDA compiler than its OpenCL one.
I have not read that paper in full and am repeating the published summary, which is exactly the sort of secondhand number this section is warning about.
The modern version of this result is more encouraging and comes from chipStar, which compiles unmodified CUDA and HIP to OpenCL and SPIR-V.
Its IWOCL 2026 keynote reported performance competitive with vendor-native toolchains across Intel discrete and integrated GPUs, AMD and NVIDIA GPUs through Rusticl, Arm Mali-G52, RISC-V systems with PowerVR graphics, and x86 and Arm CPUs, with overhead negligible on some platforms and a reasonable trade on others.
It was validated on real codes, including a quantum chemistry package with more than 20,000 lines of GPU kernels and the libCEED finite-element library.
The cheapest demonstration of the gap is the one sitting in this article. Running the three GEMM kernels above on a CPU device (PoCL’s pthread driver on a Xeon, where __local memory is ordinary RAM and there is no scratchpad to tile into), the ranking inverts completely.
Every optimisation that would win on a GPU loses here, consistently: the tiled kernel by 1.45× to 1.78×, the register-tiled one by 1.14× to 1.35×, at every size tested.
Staging a tile into __local is a copy the CPU did not need, and the barriers are pure overhead when a work-group is executed by one thread.
This is one device and it proves nothing about any GPU. What it does show is that the tuning is not incidental to the kernel: it is the kernel, and it is aimed at a memory hierarchy the target may not have.
So: OpenCL gives you functional portability for free and performance portability only if you do the work. A kernel tuned for a 64-wide wavefront and a 64 KB scratchpad will run correctly and badly on a device with a 16-wide sub-group and 32 KB, or on a CPU with neither.
The 2.x detour and the 3.0 reset
OpenCL 2.0 (18 November 2013) was an ambitious release. It added shared virtual memory, device-side enqueue, the C11-derived memory model with scoped atomics, the generic address space, pipes, program-scope global variables, and read-write images.
Every one of these was a reasonable answer to a real limitation. Together they were too much to ask of every implementer at once.
What followed is the part that shaped the next decade. AMD and Intel implemented 2.0. NVIDIA did not: it stayed on OpenCL 1.2 as its conformance level from 2015 until 2021, offering a partial 2.0 evaluation driver from February 2017 that was never conformant. Mobile vendors implemented selectively.
So a developer writing OpenCL faced a choice between targeting 1.2 and reaching everyone, or targeting 2.0 and excluding the largest installed base of discrete GPUs.
Almost everyone chose 1.2, which meant SVM, device-side enqueue, and the new memory model went largely unused, which meant there was little pressure on anyone to implement them.
OpenCL 2.1 (16 November 2015) added SPIR-V ingestion, sub-groups, clCloneKernel, low-latency device timer queries, and the OpenCL C++ kernel language. OpenCL 2.2 (16 May 2017) brought OpenCL C++ into core along with SPIR-V 1.2 and pipe storage. Both landed on an ecosystem that had not adopted 2.0, and neither changed that.
At SIGGRAPH 2017 Khronos announced that OpenCL would converge with Vulkan. This was widely reported as OpenCL being merged into Vulkan and discontinued, which was not quite what was said but was a reasonable reading of the slides.
The actual outcome was different and, in retrospect, better: the convergence happened at the IR level, through SPIR-V, and produced clspv and clvk, a compiler and runtime that let OpenCL C kernels execute on Vulkan drivers. Adobe used exactly this to ship Premiere Rush on Android. Meanwhile OpenCL kept its own roadmap under the working title “OpenCL Next”.
That became OpenCL 3.0, provisional on 27 April 2020 and final on 30 September 2020. Its central move was to invert the compatibility model: OpenCL 1.2 becomes the mandatory baseline, and every 2.x feature becomes optional and queryable.
The immediate effect was that vendors who had been stuck could ship a “current” version. NVIDIA became OpenCL 3.0 conformant on Windows and Linux with the R465 driver in April 2021, covering Maxwell and later, and exposing a specific set of optional pieces: RGBA vector component naming, the pragma unroll hint, opencl_3d_image_writes, the clCreate*WithProperties entry points, clSetContextDestructorCallback, clCloneKernel, and clEnqueueSVMMigrateMem.
Intel’s NEO compute-runtime had shipped OpenCL 3.0 for Tiger Lake from version 20.41 in October 2020, including complete 2.0 and 2.1 optional functionality and parts of 2.2.
The so called criticism of 3.0, that “everything is optional” means the standard guarantees nothing, is half right. It is true that a conformant OpenCL 3.0 device may expose exactly the OpenCL 1.2 feature set and nothing more. It is also true that this describes what the ecosystem already was; 3.0 made the situation legible instead of pretending otherwise.
The query surface it introduced is genuinely usable, and the alternative, a specification that most vendors ignore, is worse.
There are real backward-compatibility traps, though, and they are not always obvious. Program-scope global variables are an example: they were introduced in OpenCL C 2.0, and under 3.0 they are optional (__opencl_c_program_scope_global_variables).
Code that built fine against a vendor’s OpenCL 2.0 compiler can silently produce wrong results, not a build error, on the same vendor’s 3.0 driver if the driver drops the feature and the compiler’s handling of the declaration changes.
Users of PyOpenCL’s ElementwiseKernel hit precisely this on NVIDIA’s R465 driver. “OpenCL 1.2 applications run unchanged on OpenCL 3.0” is true. “OpenCL 2.x applications run unchanged” is true only if the driver still supports every 2.x feature they use, and you must check.
What OpenCL 3.1 changed
Released 4–5 May 2026; specification revision 3.1.1 dated 22 May 2026. The design philosophy is stated explicitly by the working group and is worth quoting in substance: features are proven in the field as extensions first, watched across multiple implementations, refined on developer feedback, and only then promoted to core. Everything mandated in 3.1 was already shipping somewhere.
Mandatory SPIR-V ingestion, plus mandatory support for the SPIR-V query extension. Covered in section 8; this is the change with the largest downstream effect.
Sub-groups in core, including shuffles, rotations, and an expanded set of supported data types. Applications no longer need extension guards or fallback paths for the collective operations that tuned reductions, scans, and matrix kernels are built from.
Integer dot products in core, including saturating and accumulating variants, together with extended bit operations. Both map to dedicated instructions on a wide range of modern silicon, the dp4a-class instructions, and both are the arithmetic primitives underneath INT8 inference. Promoted from cl_khr_integer_dot_product and cl_khr_extended_bit_ops.
A suggested local work-group size query in core, from cl_khr_suggested_local_work_size.
A standard device UUID query in core, matching Vulkan’s VkPhysicalDeviceIDProperties::deviceUUID. This lets an application correlate the same physical device across OpenCL and Vulkan, which is required for external memory sharing and for any sane device selection policy on a multi-GPU box. Promoted from cl_khr_device_uuid.
printf gains z and t length modifiers, for size_t and ptrdiff_t. Device-side printf could not previously format pointer-sized values without casts or format-string tricks, a small thing that shows up constantly in debugging.
CL_DEVICE_HOST_UNIFIED_MEMORY semantics clarified so it reliably distinguishes integrated from discrete GPUs.
Local memory kernel arguments may be set to zero, meaning “no local memory needed”. Kernels that opportunistically use the scratchpad no longer need a separate code path for the configuration where they don’t.
Observing CL_COMPLETE is now a synchronisation point.
The inclusive scopes rule is relaxed in the memory model, so a finer-grained scope can satisfy a coarser-grained requirement.
Implementations were in flight at release from Arm, Imagination, Intel and Qualcomm, plus Rusticl, PoCL and clvk. The first conformant listing arrived on 14 July 2026: Rusticl on Apple M1/M2 hardware under Asahi Linux, with radeonsi and Zink submissions in progress.
Mesa 26.2, released 5 August 2026, ships OpenCL 3.1 support for Rusticl on Asahi, Iris, radeonsi, llvmpipe and Zink, along with cl_khr_subgroup_rotate on radeonsi and Iris and other subgroup extensions.
The forward roadmap named by Khronos: command buffers, unified shared memory, cooperative matrix operations, low-precision AI data types including int4 and fp8, improvements to external memory sharing, and image tiling controls.
Beyond extensions, the group says it is exploring OpenCL’s role as a substrate for higher-level programming models, in safety-critical markets, and on NPUs and RISC-V accelerators.
Extensions
Extensions are named cl_khr_* (ratified, cross-vendor), cl_ext_* (multi-vendor but not ratified), or cl_<vendor>_*. A device advertises them in CL_DEVICE_EXTENSIONS as a space-separated string, or since 3.0 in the structured CL_DEVICE_EXTENSIONS_WITH_VERSION. Language-level extensions are enabled in the kernel with #pragma OPENCL EXTENSION cl_khr_fp16 : enable.
Provisional extensions carry version numbers below 1.0 and are subject to change; cl_khr_command_buffer spent years at 0.9.x. Some are shipped behind experimental headers.
That is a feature of the process: it is how features get tested before they are mandated. But it means “supports cl_khr_command_buffer“ is not a stable claim without a version.
The families worth knowing:
Numeric and language. cl_khr_fp16, cl_khr_fp64, cl_khr_int64_base_atomics, cl_khr_int64_extended_atomics, cl_khr_3d_image_writes, cl_khr_integer_dot_product, cl_khr_extended_bit_ops, cl_khr_expect_assume (compiler hints), cl_khr_kernel_clock (added provisionally in OpenCL 3.0.16, April 2024, for in-kernel profiling; developed by Arm, Imagination, Intel and Qualcomm).
Sub-groups. cl_khr_subgroups plus _extended_types, _non_uniform_vote, _ballot, _non_uniform_arithmetic, _shuffle, _shuffle_relative, _clustered_reduce, _rotate. Mostly subsumed by 3.1 core.
Execution. cl_khr_command_buffer, cl_khr_command_buffer_mutable_dispatch, cl_khr_command_buffer_multi_device, cl_khr_suggested_local_work_size, cl_khr_priority_hints, cl_khr_throttle_hints.
Memory. cl_khr_unified_svm, cl_intel_unified_shared_memory, cl_khr_extended_async_copies, cl_khr_async_copy_fence, cl_khr_device_uuid.
Interop. cl_khr_gl_sharing and cl_khr_gl_event for OpenGL; cl_khr_egl_image and cl_khr_egl_event for EGL; cl_khr_d3d10_sharing, cl_khr_d3d11_sharing, cl_khr_dx9_media_sharing on Windows; and the modern replacements: cl_khr_semaphore, cl_khr_external_semaphore with its opaque_fd and sync_fd variants, and cl_khr_external_memory with dma_buf, opaque_fd and win32 variants.
All eight were finalised together in OpenCL 3.0.16 and are the correct way to share memory and synchronisation primitives with Vulkan. NVIDIA collaborated on these specifically for Vulkan interop and ships samples using them.
Machine learning. OpenCLML on Adreno, exposed through Qualcomm’s OpenCL ML SDK. It does not appear in the Khronos registry (the cl_qcom_* extensions registered there cover host pointers, ION and Android native buffers, and performance hints), so treat it as SDK-delivered rather than as a registry extension.
This is a vendor extension that provides accelerated neural network operations, and Qualcomm reported at IWOCL 2026 that adding new accelerated ops in OpenCLML extension version 5 doubled prefill performance for their generative AI models in TVM’s Relax pipeline.
Cooperative matrix, the newest and the one that matters most for inference. On 29 April 2026 the working group published a draft cl_khr_cooperative_matrix, developed with Arm, Intel and Qualcomm. It lets an OpenCL implementation accept SPIR-V modules using SPV_KHR_cooperative_matrix, the same extension Vulkan standardised, providing cooperative load, store and multiply-add at sub-group scope, with supported matrix shapes, component types and saturation behaviours queried through a new clGetDeviceCooperativeMatrixInfoKHR.
A companion extension to expose the same capability directly in OpenCL C is in progress, with an RFC on LLVM Discourse proposing the Clang front-end changes: a cooperative matrix type attribute, built-in load/store/MAD functions, and lowering to SPIR-V-friendly LLVM IR through target extension types.
The absence of this has been a concrete, measurable limitation. Qualcomm’s IWOCL 2026 paper on llama.cpp says so plainly: unlike APIs with native cooperative matrix or tensor core abstractions, OpenCL has no standardised interface for them, so their team had to engineer portable GEMM implementations for dense and mixture-of-experts workloads by hand, using adaptive tiling, subgroup-aware parallelisation and device-specific kernel variants, to reach high utilisation without dedicated matrix instructions.
Every generation of hardware that adds matrix units widens the gap that this extension is meant to close.
The implementation landscape in 2026
NVIDIA. OpenCL 3.0 conformant since R465 (April 2021), Maxwell and later, x86/x86-64 Linux and Windows only. Functional and maintained; not where NVIDIA’s effort goes. External memory and semaphore extensions are supported for Vulkan interop.
AMD. OpenCL ships as part of CLR (Compute Language Runtimes), the repository that also contains the HIP runtime, sharing the ROCclr device layer. Development moved there from ROCm-OpenCL-Runtime at ROCm 5.6, and kernel compilation goes through the same Clang/LLVM AMDGPU path as HIP. AMD’s own effort is overwhelmingly directed at HIP, and OpenCL rides along on the shared runtime rather than being driven forward on its own; on Linux, Rusticl has at times outperformed the ROCm OpenCL stack on the same hardware.
Intel. The most complete implementation. NEO / intel/compute-runtime on Linux and Windows, OpenCL 3.0 with the full 2.0 and 2.1 optional feature set plus parts of 2.2, USM, and an unusually large set of vendor extensions. Intel also maintains the OpenCL Intercept Layer and contributes heavily to the working group: Ben Ashbaugh of Intel gave the OpenCL state-of-the-union at IWOCL 2026.
Arm. Mali GPUs, OpenCL 3.0 conformant from Mali-G78, G310, G510, G610, G710 and G78AE onward, with G615 and G715-Immortalis listed in October 2022. Arm implemented cl_ext_cxx_for_opencl and is a co-author of the cooperative matrix work.
Qualcomm. Adreno, OpenCL 3.0 full profile on current Snapdragon platforms, plus an OpenCL SDK, the OpenCL ML SDK, the Snapdragon Profiler, and a published Adreno OpenCL best-practices guide. Qualcomm is, on the evidence of the last two years, the most active silicon vendor in OpenCL: the llama.cpp backend, the TVM and MLC work, OpenCLML, and co-authorship of command buffers and cooperative matrix.
Imagination, VeriSilicon (Vivante GPU IP, OpenCL 3.0 and 1.2 full profile on the VIP9000 series for automotive and edge AI), Texas Instruments (DSP platforms), Samsung, and Cadence all appear on the Khronos conformant list.
PoCL. Portable Computing Language, MIT licensed, fifteen years old as of IWOCL 2026, developed largely at Tampere University. Conformant for CPU and Level Zero GPU targets.
PoCL 7.2-RC1 (August 2026) achieved OpenCL 3.0 conformance for CPU devices on both x86-64 (submitted on a Ryzen 9 9900X) and RISC-V (on the Star64 board and Milk-V Jupiter), and added cl_khr_extended_bit_ops, cl_khr_device_uuid, cl_khr_suggested_local_work_size, cl_khr_integer_dot_product and cl_khr_kernel_clock, with LLVM 22 support for CUDA and Level Zero back ends and LLVM 22/23 for the CPU back end.
PoCL also has a remote back end that transparently offloads OpenCL work to other machines over the network, with real memory management and distributed command scheduling rather than naive call forwarding.
Rusticl. Mesa’s OpenCL implementation on top of Gallium, written in Rust, led by Karol Herbst at Red Hat. Merged into mainline Mesa in September 2022 and shipped in Mesa 22.3, conformant with OpenCL 3.0 in November 2022 on 12th-generation Intel graphics via the Iris driver, and the first conformant OpenCL 3.1 implementation in July 2026.
It replaces Clover, which never supported images and was effectively abandoned. It requires RUSTICL_ENABLE to advertise devices for most drivers, since enabling by default is a per-driver opt-in.
Layered implementations, the strategic development of the last five years. clvk (Kévin Petit) implements the OpenCL API on Vulkan, using clspv to compile kernels; OpenCLOn12 layers OpenCL over Direct3D 12 through Mesa Gallium, which is how Windows on Arm got OpenCL; Ancle and Rusticl-over-Zink cover further combinations.
The consequence is that “does this platform have OpenCL?” increasingly means “does this platform have Vulkan or D3D12?”, and both are close to universal. chipStar’s keynote makes the point concretely: through clvk and Rusticl, CUDA code compiled to OpenCL can reach macOS GPUs by way of MoltenVK and Metal, and mobile GPUs on Android and iOS.
Conformance itself is worth explaining. The Khronos Conformance Test Suite has been open source on GitHub since 2017. Passing it and submitting the results to the Adopters Program is what entitles a product to be listed as conformant and to use the trademark; you do not need to be a Khronos member to become an adopter.
This is a meaningfully stronger regime than “we support OpenCL”: a listing on the conformant products page is a dated, versioned claim about a specific product on specific hardware.
Where OpenCL actually runs today
Mobile and edge LLM inference. llama.cpp has had an OpenCL backend since late 2024, contributed by Qualcomm and built for Adreno first. It requires sub-group support. It is tuned for Q4_0, with --pure quantisation giving the best results, and supports Q6_K and others.
For Snapdragon X2 SoCs there is a prebuilt binary kernel library covering MUL_MAT_ID with Q4_0, Q4_1, Q4_K and MXFP4, distributed through Qualcomm’s software centre, plus an on-disk compiled-program cache.
At IWOCL 2026 the Qualcomm team reported a more than fourfold prefill speedup on GPT-OSS-20B with mixture-of-experts on Snapdragon X2 Elite, from roughly 120 tokens/s to over 500, through kernel restructuring, memory-access tuning and expert load balancing.
Their roadmap names FlashAttention-style kernels, more quantisation schemes, expanded INT8 paths, and auto-tuning across vendors. That paper won the IWOCL 2026 outstanding short paper award.
Deep learning compilers. TVM and MLC have long-standing Adreno OpenCL support, and Qualcomm has upstreamed the Adreno enhancements (texture paths, specialised layouts, memory management, OpenCLML integration) into TVM’s Relax pipeline after the community deprecated Relay. Notably, they are now adding a Vulkan backend in parallel, reusing more than 90% of the target-independent optimisations, specifically to get access to Vulkan’s cooperative matmul. That is a clear signal about what the missing cooperative matrix extension costs OpenCL.
CUDA portability. chipStar compiles unmodified CUDA and HIP into fat binaries built on OpenCL and SPIR-V. Unlike source-to-source translators, it preserves the programming model and produces binaries that run without recompilation across Intel, AMD, NVIDIA, Mali, PowerVR-on-RISC-V, and CPUs.
SYCL. SYCL began as a single-source C++ layer over OpenCL. SYCL 2020 generalised to multiple backends (Level Zero, CUDA, HIP) but OpenCL remains a first-class one, and it is the backend that gives SYCL reach onto hardware without a vendor SYCL implementation. Intel’s “SYCL Everywhere” work at IWOCL 2026 described exactly this: LLVM evolving for SYCL compilation, SPIR-V as the IR, OpenCL as the portability layer for accelerator offload, and PoCL and Mesa quietly extending device coverage. There is a nice concrete instance of the dependency in the other direction: the FunGT graphics engine uses SYCL’s OpenCL backend specifically to reach cl_khr_gl_sharing, because SYCL has no native OpenGL interop.
FPGAs. Intel’s FPGA SDK for OpenCL and, historically, Xilinx SDAccel (Khronos-conformant since January 2015, later folded into Vitis) used OpenCL as an HLS front end. Pipes and pipe storage exist in the specification largely because of this constituency. The centre of gravity here has moved toward vendor HLS flows and oneAPI, but OpenCL-derived tooling persists in shipping products.
Safety-critical and the road not taken. Khronos says it is exploring OpenCL’s role in safety-critical markets, and there is precedent in the neighbourhood: Vulkan SC exists, SYCL SC is in development, and IWOCL 2026 carried a paper on functional-safety-oriented GPU development that migrates CUDA to SYCL and then applies static analysis to strip constructs unsafe under IEC 61508 and ISO 26262. There is no OpenCL SC. Whether one appears is a reasonable proxy for how seriously the automotive and industrial constituency takes the standard.
WebCL is the road that was not taken. Khronos formed the working group in March 2011 and published a 1.0 specification in March 2014, aiming to expose OpenCL to JavaScript. No browser shipped it; the security surface of arbitrary compute kernels in a web page proved unattractive, and Khronos now lists it among its inactive standards. The niche it aimed at is being filled by WebGPU instead.
The long tail. darktable, GIMP, Blender’s historical Cycles OpenCL backend, LuxMark, Leela Chess Zero, LAMMPS, GROMACS, ffmpeg filters, Arm Compute Library, ncnn, MNN, and the TensorFlow Lite GPU delegate. This is not glamorous work and it is a great deal of deployed code.
Why CUDA really won
It is worth being accurate about this rather than reaching for “NVIDIA had better marketing.”
Single-source versus split-source. In CUDA, host and device code live in the same translation unit, compiled by one compiler that type-checks kernel launches. In OpenCL, kernels are strings compiled by a different compiler at runtime, with arguments bound by index and size. The gap is not aesthetic; it is a difference in how many errors the type system catches. Templates work across the boundary in CUDA and did not in OpenCL until C++ for OpenCL, which arrived in 2020 and is not ratified.
Libraries. cuBLAS, cuDNN, cuFFT, cuSPARSE, Thrust, CUB, NCCL: vendor-maintained, aggressively tuned, and free. The OpenCL equivalents were clBLAS (deprecated), CLBlast (excellent, and maintained by a far smaller group), clFFT, VexCL, ArrayFire, Boost.Compute. Building a GPU application on CUDA meant assembling tuned components; on OpenCL it frequently meant writing them.
Asymmetric investment, for structural reasons. NVIDIA’s compiler team optimises one language for one family of architectures. An OpenCL implementation is a compiler plus a runtime that must be correct across a specification designed to accommodate CPUs, DSPs, FPGAs and fixed-function accelerators. Vendors were also, understandably, not motivated to make the vendor-neutral API as fast as their proprietary one. A compiler team’s attention is a budget, and it goes where the strategic return is.
The 2.x fragmentation. Covered above. Between 2013 and 2020 the answer to “which OpenCL version can I target?” was 1.2, which meant OpenCL was frozen at a 2011 feature set during precisely the years when GPU compute exploded.
Tooling. Nsight Compute and Nsight Systems have no OpenCL equivalent. There are good OpenCL tools, the Intercept Layer, clinfo, clpeak, Oclgrind for memory error detection, Snapdragon Profiler, Intel VTune, now CLVizulayer, but they are assembled from several projects rather than shipped as one supported product.
Apple’s exit. Apple invented OpenCL, held the trademark, and deprecated it in macOS 10.14 Mojave in 2018 in favour of Metal, along with OpenGL, telling developers to move computational work to Metal Performance Shaders. Losing your originator is a bad look regardless of installed base.
Blender’s retirement of the OpenCL path in Cycles is the compact version of the whole story. The stated reasons were a limited kernel implementation, driver bugs, and a stalled standard, and the combination made maintenance untenable. Not one of those is a flaw in the specification. All three are consequences of the specification being implemented by people who were investing elsewhere.
What is also true, and less often said: a striking amount of what OpenCL specified early has since become the industry consensus. SPIR-V is now the IR for Vulkan, OpenCL, and SYCL, and is a target for Slang, clspv, and a growing set of DSLs. Sub-groups are warps by another name and every API now exposes them.
Command buffers are CUDA Graphs and Level Zero command lists. Unified shared memory is CUDA’s unified memory with a better capability model. Cooperative matrix is the same abstraction in Vulkan and OpenCL, standardised jointly. The ideas held up; what OpenCL never got was the distribution to go with them.
Writing OpenCL in 2026
Practical positions, offered as opinion rather than doctrine:
Target OpenCL 1.2 plus queried features, or 3.0 with capability checks. Do not target 2.x. Query at startup, build the kernel variant the device can run, and keep the fallback paths honest by testing them: PoCL and Rusticl on a laptop will exercise a very different feature set from a vendor GPU driver.
Ship SPIR-V, once 3.1 drivers are common. Until then, ship source with a persistent binary cache keyed on device name, driver version, source hash and build options. The startup compile cost is real and users notice it.
Concretely, the startup path that makes all of this work is short enough to paste:
/* Select a device, discover what it can actually do, and build the
matching kernel variant. This is the whole portability story. */
cl_device_id dev = pick_device();
char prof[64], ver[128];
clGetDeviceInfo(dev, CL_DEVICE_PROFILE, sizeof(prof), prof, NULL);
clGetDeviceInfo(dev, CL_DEVICE_VERSION, sizeof(ver), ver, NULL);
int embedded = (strcmp(prof, "EMBEDDED_PROFILE") == 0);
int major = ver[7] - '0'; /* "OpenCL M.m ..." */
int minor = ver[9] - '0';
/* OpenCL 3.0+ optionality queries; on 1.2 they simply fail, which is
the answer. Treat a failed query as "unsupported", never as fatal. */
cl_device_atomic_capabilities atomics = 0;
cl_bool generic_as = CL_FALSE, images = CL_FALSE;
cl_device_device_enqueue_capabilities dq = 0;
clGetDeviceInfo(dev, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
sizeof(atomics), &atomics, NULL);
clGetDeviceInfo(dev, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT,
sizeof(generic_as), &generic_as, NULL);
clGetDeviceInfo(dev, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES,
sizeof(dq), &dq, NULL);
clGetDeviceInfo(dev, CL_DEVICE_IMAGE_SUPPORT, sizeof(images), &images, NULL);
size_t next; clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, 0, NULL, &next);
char *ext = malloc(next);
clGetDeviceInfo(dev, CL_DEVICE_EXTENSIONS, next, ext, NULL);
int subgroups = strstr(ext, "cl_khr_subgroups") != NULL || (major > 3)
|| (major == 3 && minor >= 1);
int fp16 = strstr(ext, "cl_khr_fp16") != NULL;
int fp64 = strstr(ext, "cl_khr_fp64") != NULL;
/* Build options carry the decisions into the kernel. */
char opts[512];
snprintf(opts, sizeof opts,
"-cl-std=%s -DSUBGROUPS=%d -DFP16=%d -DFP64=%d -DIMAGES=%d "
"-DEMBEDDED=%d -cl-kernel-arg-info",
(major >= 3) ? "CL3.0" : "CL1.2",
subgroups, fp16, fp64, images == CL_TRUE, embedded);
cl_program p = build_from_cache_or_source(ctx, dev, opts); /* see §8 */
and the kernel side picks its own path from the same flags, preferring the mechanism the device actually has:
inline float block_sum(float v, __local float *scratch) {
#if SUBGROUPS && (defined(__opencl_c_subgroups) || defined(cl_khr_subgroups))
# if defined(cl_khr_subgroups) && !defined(__opencl_c_subgroups)
# pragma OPENCL EXTENSION cl_khr_subgroups : enable /* needed on 2.x */
# endif
return sub_group_reduce_add(v);
#elif defined(__opencl_c_work_group_collective_functions)
return work_group_reduce_add(v);
#else
size_t l = get_local_id(0), n = get_local_size(0);
scratch[l] = v;
barrier(CLK_LOCAL_MEM_FENCE);
for (size_t s = n >> 1; s > 0; s >>= 1) {
if (l < s) scratch[l] += scratch[l + s];
barrier(CLK_LOCAL_MEM_FENCE);
}
return scratch[0];
#endif
}
Three variants of one reduction is not elegant, and it is the actual cost of the portability the standard sells. The fallback path is the one to test hardest, because it is the one that runs on the hardware you did not anticipate.
Use -cl-kernel-arg-info, set a context error callback, and check CL_KERNEL_PRIVATE_MEM_SIZE. These three cost nothing and catch a disproportionate share of problems.
Treat the local work size as a measured parameter, not a guess or a default. A 6.8× spread across plausible values is normal. Benchmark NULL against a handful of multiples of CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, bounded by CL_KERNEL_WORK_GROUP_SIZE, or use the 3.1 suggested-size query. Sometimes the runtime wins; assume neither.
Assume nothing about work-group execution order. No inter-work-group spin-waiting, ever.
Treat tuning parameters as data, not code. Tile sizes, work-per-thread, vector width and unroll factors should be build-time defines selected from a per-architecture table, with autotuning as the fallback. This is what CLBlast does and it is the only approach that survives contact with five vendors.
Reach for the ecosystem before writing kernels. CLBlast for BLAS, VkFFT or clFFT for transforms, Arm Compute Library on Mali, OpenCLML on Adreno. And if the actual goal is portable C++ rather than portable kernels, SYCL over an OpenCL backend is very likely the better answer: OpenCL is a good compilation target and a mediocre application-level programming model, and the working group’s own framing now treats it that way.
Debug with the layer mechanism. The Intercept Layer will tell you what your application is really doing to the driver, without a rebuild.
Where this goes
The OpenCL of 2026 has settled into a role nobody planned for it in 2008. It is not the API most application developers write. It is the layer that other things stand on: SYCL implementations, chipStar, TVM, MLC, llama.cpp’s mobile path, the Mesa and PoCL stacks that give OpenCL, and therefore everything above it, reach onto hardware whose vendors have shipped nothing.
The working group’s own language has shifted to match, talking about OpenCL as a substrate for higher-level programming models rather than as the model itself.
The most urgent open question is cooperative matrix. Hardware matrix units are where the FLOPs are, Vulkan standardised access to them years ago, and until the OpenCL extension is final and shipping, every inference backend on OpenCL is hand-rolling GEMM against silicon it cannot fully address.
Qualcomm adding a Vulkan path to TVM alongside the OpenCL one, explicitly to get cooperative matmul, is the warning shot. The low-precision data types are the same problem one layer down: int4 and fp8 are on the roadmap, and the workloads that need cooperative matrix need those too.
Underneath both sits the question of whether the 3.1 SPIR-V mandate converts into broad driver support quickly enough to matter to the compiler authors it was written for.
The encouraging signal is that the working group has stopped repeating the 2.0 mistake. Everything mandated in 3.1 was already deployed. The extension pipeline is public, the drafts are on GitHub, the Clang RFCs are on Discourse, and the conformance suite is open source.
Whether that process moves fast enough against a competitor with a decade’s head start and no committee is a genuinely open question, and I do not think anyone should be confident either way.
What I’d say with more confidence is that the obituaries were premature by about a decade, and that if you are building anything that has to run on hardware you do not control (phones, embedded silicon, whatever a customer happens to own), OpenCL is currently the only thing in the category that is both open and actually there.
Which returns to where this started. A standard that has to describe hardware it cannot see can only ever specify floors. The measurements in this article are all instances of that: an accuracy bound four orders of magnitude looser than what an implementation delivers, a work-group size the specification will not choose for you and that costs 6.8× to get wrong, an allocation flag that means nothing until you change how you use it.
Every one of those is the standard declining to promise something it cannot promise for every device.
That is the failure mode and it is also the whole value. CUDA can tell you what your hardware does because NVIDIA built it. OpenCL cannot, and in exchange it is the reason a kernel written in 2011 runs today on a phone, an FPGA, a RISC-V board, and an Apple GPU under a driver Apple did not write and does not support.
The interesting question was never whether that trade was worth it in general.
It is whether it is worth it for the thing you are building, and the honest answer is that for most people writing CUDA in 2026 it is not, and for the people shipping software onto hardware they will never see it is the only trade on offer.
Appendix A: version timeline
Appendix B: queries worth calling at startup
/* Identity and version */
CL_PLATFORM_VERSION CL_DEVICE_VERSION
CL_DEVICE_NAME CL_DEVICE_VENDOR
CL_DRIVER_VERSION CL_DEVICE_OPENCL_C_ALL_VERSIONS /* 3.0 */
CL_DEVICE_UUID_KHR /* core in 3.1 */
/* Shape of the machine */
CL_DEVICE_MAX_COMPUTE_UNITS CL_DEVICE_MAX_WORK_GROUP_SIZE
CL_DEVICE_MAX_WORK_ITEM_SIZES CL_DEVICE_LOCAL_MEM_SIZE
CL_DEVICE_LOCAL_MEM_TYPE CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
CL_DEVICE_MAX_MEM_ALLOC_SIZE CL_DEVICE_GLOBAL_MEM_SIZE
CL_DEVICE_HOST_UNIFIED_MEMORY /* meaningful as of 3.1 */
/* Capabilities */
CL_DEVICE_EXTENSIONS_WITH_VERSION
CL_DEVICE_SVM_CAPABILITIES
CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES
CL_DEVICE_ATOMIC_FENCE_CAPABILITIES
CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES
CL_DEVICE_PIPE_SUPPORT
CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT
CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT
CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT
CL_DEVICE_IMAGE_SUPPORT
CL_DEVICE_SINGLE_FP_CONFIG CL_DEVICE_DOUBLE_FP_CONFIG
/* Per kernel, after build */
CL_KERNEL_WORK_GROUP_SIZE
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
CL_KERNEL_LOCAL_MEM_SIZE
CL_KERNEL_PRIVATE_MEM_SIZE
CL_KERNEL_COMPILE_WORK_GROUP_SIZE
Appendix C: the same concept in seven APIs
Every modern compute API converged on the same execution model with different nouns. This is the translation table, and it is most of what porting between them consists of.
Two structural observations fall out of the table. The rows are almost perfectly aligned, which is why source-to-source translation between these APIs works as well as it does and why chipStar can compile CUDA onto OpenCL at all. And the one row where OpenCL is behind, the last one, is the row where the FLOPs are.
The source model row is the difference that decided the market. Single-source means host and device code are compiled together by one compiler, so a kernel launch is type-checked and templates cross the boundary.
Split-source means kernels are strings or binaries bound by index at runtime. The pattern is not simply proprietary-versus-open, and Metal is the counterexample that shows it: Apple controls its compiler completely and still chose split-source, because MSL is its own language in its own files.
What actually predicts the choice is whether one compiler is expected to consume both halves. CUDA, HIP and SYCL say yes and get type-checked launches and templates across the boundary.
OpenCL, Vulkan, Metal and Level Zero say no, and pay for it at every call site where an argument is bound by index instead of by name.
Appendix D: sources
Khronos OpenCL Registry and the unified API, C, and SPIR-V Environment specifications (rev. 3.1.1, 22 May 2026); the Khronos blog posts announcing OpenCL 3.1 (4 May 2026) and the cooperative matrix extensions (29 April 2026);
the OpenCL-Docs, OpenCL-Headers, OpenCL-ICD-Loader and OpenCL-CTS repositories; the IWOCL 2026 program and papers, particularly the Qualcomm llama.cpp paper, the chipStar keynote, Intel’s AI-workload talk, and the CLVizulayer presentation;
NVIDIA’s OpenCL 3.0 conformance announcement (April 2021); llama.cpp’s OpenCL backend documentation; PoCL and Mesa release notes; Clang’s OpenCL support documentation; and Phoronix’s reporting on the 3.0.x and 3.1 releases and on conformance submissions.