The Contract a Compute API Encodes

A GPU cannot be programmed by calling a function. Its execution units are fed by a command processor that consumes buffers of encoded commands, its memory is a separate pool reachable through a separate address space, and a kernel has no call frame that a host thread could push and pop. Every acceleration interface, whether tied to one vendor or designed to span several, must therefore answer the same four questions before it can run anything: what is the code, how is the index space shaped, which memory does the code touch, and in what order do operations happen.

The differences between CUDA, OpenCL, Vulkan compute, SYCL, and OpenMP offload are largely differences in how those four answers are spelled, not differences in what has to be answered. Reading an unfamiliar API reduces to locating its answer to each question.

Code that is compiled for a device and entered without a call

A kernel is a function that the host cannot call. It is compiled by a device compiler, emitted into a binary the device can load, and entered by the command processor with an index-space descriptor rather than an argument list in the ordinary sense. The parameters do arrive as an argument buffer: CUDA passes them by value in a small constant bank, OpenCL packs them into the kernel’s argument slots, Vulkan passes a push-constant or storage-buffer block. There is no return value and no call stack in the host sense.

That single constraint forces the rest of the design. Since the host cannot call the function, the host must instead describe a region of the index space over which the function will be instantiated. Since there is no return path, results must be written into memory the host can later read. And since the function cannot allocate, all working storage has to be declared before launch, either in the kernel’s own scratch space or in an object the host created.

The code itself is not portable machine code. It is either a virtual instruction set that a device driver lowers at load time (PTX for NVIDIA, SPIR-V for Vulkan and OpenCL 2.x), or native code for one specific architecture. The distinction between the two is the subject of PTX, SASS, and What a Fat Binary Contains, and it decides whether a binary written today runs on hardware released later.

The index space is a description, not a thread list

Nowhere in a launch is there a list of threads. The host supplies extents, and the hardware derives the rest. CUDA names them a grid of blocks and a block of threads; OpenCL names them an NDRange of work-groups and work-items; Vulkan names them a workgroup count and a local size; SYCL names them an nd_range. All four denote the same two-level structure, and in all four the two levels differ in kind rather than degree.

A work-item is a lane in a SIMT machine: it shares an instruction stream with its neighbours in a group of 32 on NVIDIA hardware, 64 on older AMD designs, and a size the implementation chooses elsewhere. A work-group is the unit that can synchronise, because its members are guaranteed to be resident on one multiprocessor at the same time. This is why the barrier primitive takes effect only within a group, whether it is __syncthreads() in CUDA, barrier(CLK_LOCAL_MEM_FENCE) in OpenCL, or workgroupBarrier() in WGSL. None of these primitives is device-wide, so a kernel that needs to synchronise the whole grid must either be split into two launches or, in CUDA, use the cooperative-groups grid synchronisation, which requires every block of the grid to be co-resident.

The same three-way split, spelled in two of the interfaces:

// CUDA: grid of blocks, block of threads, both three-dimensional.
kernel<<<dim3(gx, gy, 1), dim3(256, 1, 1)>>>(ptr, n);

// OpenCL: one work dimension, global extent, local extent, offset.
size_t global[1] = { gx * 256 };
size_t local[1]  = { 256 };
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, local, 0, NULL, NULL);

The CUDA form declares the grid extent directly and derives the thread count from the block shape. The OpenCL form declares the total extent and the local extent separately, so the number of groups is a quotient the runtime computes, and a global extent that is not a multiple of the local one is rejected rather than rounded.

The nesting also determines where a work-item’s identifier comes from. A three-dimensional index is flattened against the declared extents, and the arithmetic of that flattening is what Launch Geometry and the Arithmetic of Thread Indexing works through. The important point here is that the API exposes the shape, not the mapping: nothing in the contract promises that consecutive identifiers land on consecutive hardware lanes, and on most hardware they do so only when the local size is a multiple of the group width.

Memory objects, and the address spaces they name

A compute API has to give names to at least three kinds of storage, and the naming is where portability breaks first.

Device global memory is the large pool. It is allocated through the API (cudaMalloc, clCreateBuffer, vkAllocateMemory, sycl::malloc_device), and its lifetime is managed by the host, not by the kernel. Kernel-visible local scratch is the small per-group region: CUDA’s __shared__, OpenCL’s __local, SPIR-V’s Workgroup storage class, typically tens of kilobytes per multiprocessor. Registers are the third and smallest, allocated by the compiler and effectively an allocation the host influences only indirectly, through the block size it requests.

The three have sharply different access costs, and the difference is architectural rather than incidental. Global memory traffic goes through the memory hierarchy and is measured in hundreds of cycles when it misses; scratch memory is banked SRAM on the multiprocessor, on the order of tens of cycles; registers are effectively free. An API that hides the distinction, with a single pointer type and no address-space qualifier, forces the compiler either to infer the space conservatively or to emit generic accesses that resolve at runtime. This is the reason OpenCL C requires __global, __local, __constant, and __private qualifiers on pointer parameters, and the reason SPIR-V encodes the storage class in the pointer type itself.

Coherence is the second trap. Device memory is not cache-coherent with the host on a discrete GPU. Writes made by a kernel are visible to host reads only after an explicit transfer or a fence with system scope, and host writes are visible to a kernel only after the transfer that carried them completes. Unified memory relaxes the mechanics, since the driver migrates pages instead of copying them, but not the requirement for a synchronising operation at the boundary. The mechanisms are covered in Host-Device Transfers and the Cost of Crossing the Bus.

Queues decide what may overlap

The fourth obligation is ordering. All of these APIs are asynchronous by construction: the host enqueues work and continues, because the whole point of the device is to run while the host does something else. That raises the question of what the enqueue order means.

The default in CUDA streams and OpenCL command queues is in-order execution: operations in one queue complete in the order they were enqueued, and a kernel launched after a copy sees that copy’s results. OpenCL permits an out-of-order queue, created with CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, in which the only ordering is what event dependencies impose. Vulkan takes the opposite default from both: commands in a queue may overlap unless barriers or event waits say otherwise, and the program is responsible for declaring every hazard.

The practical consequence is that a program’s correctness depends on which of these regimes it was written against, and the differences are not cosmetic. Porting code that relied on implicit in-order semantics into a queue with explicit barriers usually produces a program that runs and produces wrong answers under load. Command Queues, Events, and Cross-Stream Ordering treats the dependency graph in detail.

What the contract leaves to the implementation

Between the four answers there is a wide band of behaviour the API does not fix. Work-group size is bounded but not prescribed: CUDA caps a block at 1024 threads since compute capability 2.0, while OpenCL’s guaranteed floor for CL_DEVICE_MAX_WORK_GROUP_SIZE is far lower, so a portable program queries the device rather than assuming a number. Scheduling order within a group is unspecified. Whether two work-items in the same group run in lockstep is unspecified, and since Volta NVIDIA hardware can diverge and reconverge independent threads within a warp, code that relied on implicit lockstep was always relying on something the contract never promised.

The pattern to read for, in any compute API, is which of the four answers is fixed by the specification and which is a query against the device. Code, index space, and queue semantics are fixed; memory sizes, group sizes, and the available instruction set are queries. Programs that hardcode the queried values are the ones that fail on the next architecture, which is why Capability Negotiation Across Architectures is a distinct problem from merely choosing an API.