Portable Single-Source Layers and Where They Leak

The vendor APIs are not portable in source, in binary, or in performance, and the layers built on top of them exist to recover the first of the three. A single-source model puts host code and kernel code in one translation unit, marks the kernel with an annotation or a lambda, and leaves the separation of compilation to the toolchain. The pattern recurs across SYCL’s parallel_for, OpenMP’s target teams distribute, and Kokkos’s parallel_for range policy: a lexical region is designated as device code, and the compiler emits two versions of it.

None of these layers can make a kernel portable in performance. What they can do is make the source portable, and the interesting engineering is in the places where they fail to do even that.

How one source file becomes two programs

The mechanism is the same in all three, and it is a compiler pass rather than a runtime feature. The compiler parses the whole translation unit, identifies the regions marked for offload, and clones them into a device-side compilation. In SYCL the clone is triggered by the device-callable requirement on a lambda passed to a submit; in OpenMP by the declare target set, which includes any function reachable from a target region; in Kokkos by the KOKKOS_LAMBDA macro, which expands to an annotation the compiler recognises.

The cloned code is compiled with a device compiler, a different front end with a restricted C++ dialect, into the target IR: SPIR-V for a generic path, native code when the toolchain supports a specific architecture. The host side keeps a stub that packages the arguments and launches the device version through the underlying vendor API. The result is that SYCL lowers to OpenCL or to a vendor backend, OpenMP offload lowers to the same vendor APIs the C interface would call, and the runtime underneath is unchanged.

This is the crucial structural fact: the single-source layers are front ends. Every limitation of the layer’s device dialect is an artefact of what the backend compiler can accept, and every performance characteristic is that of the backend, with the layer’s launch decision inserted.

The device dialect is narrower than C++

The restrictions are not stylistic. Device code has no dynamic allocation that the host can observe, no exceptions in most implementations, no virtual dispatch that would require a vtable the device can resolve, and no standard library beyond what the implementation has ported.

RTTI and exceptions are the sharpest cases. A kernel that throws cannot propagate the exception to the host, because there is no unwinding across the launch boundary; the standard-conforming response is to reject the code at compile time. dynamic_cast requires type information the device image does not carry. Where these compile, they often compile because the implementation silently permits them and the behaviour at the boundary is unspecified.

The restricted dialect extends to pointer provenance. A plain pointer captured by a device lambda carries no record of which memory it addresses, so unless the implementation can trace its origin the compiler must emit a generic dereference that resolves the address space at run time. Code written directly against the vendor API, with a __global-qualified pointer, gets a specialised load instead. The layer’s generic form is not wrong; it is one address-space resolution slower, and the difference appears in the generated code rather than in the source.

// Two ways to reach memory from a kernel body.
queue.submit([&](sycl::handler &h) {
    auto acc = buf.get_access<sycl::access_mode::read_write>(h);  // mode is in the type
    h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> i) {
        acc[i] *= 2.0f;                             // specialised load and store
    });
});

// ptr came from sycl::malloc_device, so it is device memory the type does not say.
queue.submit([&](sycl::handler &h) {
    h.depends_on(previous);                         // ordering now stated by hand
    h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> i) {
        ptr[i] *= 2.0f;                             // generic address space
    });
});

The accessor exists for exactly this reason. The buffer-accessor model that SYCL inherited from OpenCL encodes the address space and the access mode in the type, which lets the compiler generate a specialised load and lets the runtime derive dependencies without explicit events. The unified shared memory model that SYCL also offers trades that typing back for pointer convenience, and with it the automatic dependency: a kernel on a raw pointer has to declare its ordering through an explicit dependency or the queue’s own semantics, the same trade as mapped host memory against explicit transfers.

Four leaks, in increasing order of severity

Compilation leaks. A construct accepted for one backend is rejected by another. Forward declarations, nested classes in device code, constexpr evaluation depth, and the exact set of diagnostic pragmas differ. The practical consequence is that a codebase must be built against every target it claims to support, in continuous integration, or the leak is discovered by whoever tries the unsupported backend.

Address-space leaks. The same kernel performs differently depending on whether the pointer’s provenance is statically known. A layer that types its memory objects well avoids the leak; a layer that reduces everything to raw pointers exposes it. The symptom is a factor-level slowdown on one backend and none on another, with no source difference to explain it.

Launch-configuration leaks. A work-group size that is optimal or even legal on one device is invalid on another. The maximum group size is a queried property, as discussed in Capability Negotiation Across Architectures, and a layer that accepts a requested size and forwards it to a backend whose limit is lower will fail at launch. Layers that accept a range and an optional tile size, and clamp it against a query, avoid this; layers that require an explicit local size do not.

Semantic leaks. These are the expensive ones because they produce correct-looking wrong results. Reduction order is unspecified in every one of these models, so a floating-point sum over a large array differs in the low bits across backends and across work-group sizes on the same backend. A program that compares against a reference computed once, on one device, will fail on another for reasons that are entirely conforming. Atomic operations have the same problem in a different form: the memory ordering they provide is the weakest the model allows, so a lock-free algorithm that is correct under a stronger ordering is racy under the layer’s.

Cost of the front end at launch time

A second-order cost sits in the layer’s own runtime rather than in the device compiler. The launch path in a SYCL implementation goes through handler bookkeeping, accessor dependency extraction, and a command-group construction stage before it reaches the backend’s enqueue. OpenMP offload has its own set of runtime calls to marshal the argument list into a contiguous block and look up the target region.

The per-launch overhead is host-side and on the order of microseconds, against a backend enqueue that is itself of the same order. For a kernel that runs for milliseconds the difference is invisible. For a workload of many short kernels the front end can become the dominant term, and the usual workaround is the same one that applies to the raw APIs: launch fewer, larger kernels rather than many small ones.

The bookkeeping also imposes a lifetime discipline. In the buffer-accessor model, an accessor’s destruction is what completes the dependency the runtime inferred; a program that keeps accessors alive longer than needed serialises work that the runtime would otherwise overlap. In the unified shared memory model the same ordering is expressed with explicit dependency objects, and omitting them produces exactly the unsynchronised queue behaviour described in Command Queues, Events, and Cross-Stream Ordering.

What the layers are actually worth

They are worth the source portability, and the condition for collecting that value is that the abstraction boundary is drawn where the semantics are genuinely the same: index space, memory objects, launches, and the queue structure. Those are the four obligations identified in The Contract a Compute API Encodes, and all of them have the same meaning across backends.

Below that line, the layers cannot help, and code that needs the difference (a specific instruction, a specific cache configuration, a specific asynchronous copy engine) has to reach through to the backend. SYCL’s interoperability handles, OpenMP’s use_device_ptr, and Kokkos’s backend-specific memory spaces are all escape hatches for exactly this purpose. A portable codebase that never uses one is either not optimised or not portable in the sense that matters.

The pattern that survives is layered: a portable outer structure, with per-backend specialisations behind a single dispatch point, selected by capability queries and validated by tests that run on each target. That is more work than either extreme, and it is the only arrangement that keeps both properties at once.