How Is This Not SYCL?
Show anyone a new approach to heterogeneous programming and you get the same question inside thirty seconds. Sometimes it is “how is this not SYCL,” sometimes “how is this not Kokkos,” and if the person is over forty, “how is this not OpenCL.”
It is a fair question, and most answers to it are bad. The bad answer is a strawman: SYCL is verbose, SYCL is Intel-flavoured, SYCL needs a special compiler. None of that is a serious objection, and anyone who has actually shipped a SYCL codebase will know within a sentence that you have not.
The answer is worth getting right, because SYCL solved a real problem and solved it well. So this post is the tour I wish existed: what single-source heterogeneous C++ actually puts in the type system, where it stops, and — the part that took me a while to appreciate — why it stops there, which turns out to be a C++ constraint rather than a failure of imagination on anyone’s part.
I will spend most of the post on the two systems that make the strongest case: SYCL and Kokkos. Kokkos is the harder objection of the two, and almost nobody raises it first.
What SYCL actually solved
Start with what the world looked like before it.
CUDA is a language extension. __global__, __device__, __shared__, and kernel<<<grid, block>>>(args) are not C++; they are syntax nvcc understands and your other compilers do not. That has consequences beyond aesthetics. Your kernels are not ordinary functions, your build needs a second compiler that owns the whole translation unit, and the code is structurally locked to one vendor.
OpenCL went the other way. Kernels are strings, compiled at runtime by the driver, passed arguments positionally by index:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
// 1. The kernel is a string. The C++ compiler never sees it, so it cannot
// type-check anything inside it against the host code that calls it.
const char *src = "__kernel void saxpy(float a, __global float *x,"
" __global float *y) {"
" int i = get_global_id(0);"
" y[i] = a * x[i] + y[i];"
"}";
// 2. Arguments are set by position and size. Pass the wrong index, the wrong
// type, or the wrong size and nothing complains until run time -- and often
// not even then, just a wrong answer.
clSetKernelArg(kernel, 0, sizeof(float), &a);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &x_buf);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &y_buf);
Every type error that C++ would have caught in the host language is a runtime error again, because the kernel and its caller are compiled by different compilers at different times and never meet.
SYCL’s contribution is that both problems go away without inventing a dialect. A kernel is a lambda. It is ordinary C++, in the same translation unit as the host code, type-checked against its callers by the compiler you already use:
1
2
3
4
5
6
7
8
9
// 1. A plain queue. Ordinary object, ordinary constructor.
sycl::queue q;
// 2. The kernel body is a C++ lambda. `a`, `x` and `y` are captured by the
// normal rules, so a type error here is a compile error in the host TU --
// not a runtime surprise from a driver-compiled string.
q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) {
y[i] = a * x[i] + y[i];
}).wait();
No new keywords. No <<<>>>. Templates work, so do lambdas, overloads, and the parts of the standard library the device backend supports. It compiles under a conforming compiler such as Intel’s DPC++ or AdaptiveCpp, targeting CUDA, HIP, Level Zero, or plain OpenMP on the host.1
The mechanism behind that is a genuinely clever bit of engineering, and it is where “single source” gets its meaning:
graph LR
A["one .cpp<br/>host + kernels"] --> B["device compiler pass"]
A --> C["host compiler pass"]
B --> D["SPIR-V / PTX / GCN<br/>device image"]
C --> E["host object<br/>+ integration header"]
D --> F["fat binary"]
E --> F
F --> G["runtime: pick a device,<br/>load the matching image"]
style A fill:#1e3a5f,color:#fff
style G fill:#7f1d1d,color:#fff
One source file, compiled twice. The device pass extracts the kernel lambdas and emits device images; the host pass emits ordinary objects plus an integration header that lets the runtime connect a lambda’s type to its compiled image. The result is a fat binary that carries device code for every target you asked for.
That is not a small achievement, and it is the thing people mean when they say a new language “is just SYCL.” If your only claim is single-source heterogeneous compilation, they are right, and SYCL got there first with a standard behind it.
SYCL does put address spaces in the type
Here is the concession most pitches skip, and skipping it is how you lose the room.
SYCL has address spaces in its type system. It is not an oversight that someone is about to fix; it has been there from the start, inherited from OpenCL C. From the DPC++ headers:
1
2
3
4
5
6
7
8
9
enum class address_space : int {
private_space = 0,
global_space = 1,
constant_space __SYCL2020_DEPRECATED("sycl::access::address_space::constant_"
"space is deprecated since SYCL 2020") =
2,
local_space = 3,
generic_space = 4,
};
These are real type parameters2, carried by multi_ptr<T, address_space, access::decorated>, and Intel extends the set further with device_ptr and host_ptr as subsets of the global space.3 A pointer’s memory region can be part of its type, statically, today.
Two things temper it.
The first is what those spaces describe. private, local, global, constant are the regions within one device — registers, scratchpad, DRAM, read-only memory. They are the OpenCL memory model, which was designed for a single accelerator with a partitioned address space. They say nothing about which accelerator.
The second is the direction of travel. SYCL 2020 deprecated constant_space, added generic_space, and made explicit qualification largely unnecessary — plain pointers work and the compiler infers. Intel’s own extension documentation is candid that the motivation is optimisation rather than safety: annotating a USM pointer’s space lets the compiler “perform better alias analysis, which typically leads to better throughput and smaller silicon area.”3 Address spaces survived as a performance hint. The safety property they could have carried was not the point.
Where the type stops: which device
Now the part that matters.
In SYCL, a device is a runtime value. Selection happens by scoring functions, and the declarations say it plainly:
1
2
3
4
__SYCL_EXPORT int default_selector_v(const device &dev);
__SYCL_EXPORT int gpu_selector_v(const device &dev);
__SYCL_EXPORT int cpu_selector_v(const device &dev);
__SYCL_EXPORT int accelerator_selector_v(const device &dev);
Each takes a device and returns an integer score.4 The runtime enumerates what it found, scores each candidate, and keeps the winner. sycl::queue q(sycl::gpu_selector_v) does not mean “this queue is a GPU queue” in any sense the type system knows. It means “at construction, prefer whatever scored highest among things that looked like GPUs,” and if nothing qualifies you get an exception.
The allocation side matches. From usm.hpp:
1
2
__SYCL_EXPORT void *malloc_device(size_t size, const queue &q, ...);
template <typename T> T *malloc_device(size_t Count, const queue &Q, ...);
The typed overload gives back a T*.5 Not a device_ptr<T, Q>, not a pointer parameterised by the queue or device it belongs to — a T*. Allocate on two GPUs and you hold two float* values of identical type, pointing into two disjoint physical memories, freely assignable to each other and to a host float*. Dereference the wrong one in the wrong place and you get undefined behaviour, or a segfault, or silence and a wrong number.
Here is the bug that falls out, written the way it actually appears — not as a contrived example, but as the ordinary shape of multi-GPU code:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
// 1. Two queues on two different physical GPUs.
sycl::queue q0(sycl::gpu_selector_v);
sycl::queue q1(sycl::gpu_selector_v); // may or may not be a different device
// 2. Two device allocations. Note the declared types: both are `float *`.
// Nothing distinguishes them. They are assignable to one another, comparable
// to one another, and passable to any function expecting a float pointer.
float *a = sycl::malloc_device<float>(n, q0);
float *b = sycl::malloc_device<float>(n, q1);
// 3. A kernel submitted to q0, reading `b` -- which lives in q1's memory.
// This compiles. There is no diagnostic. On a system without peer access
// it is undefined behaviour; with peer access it silently runs slowly
// across the interconnect, which is worse, because it works.
q0.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) {
a[i] += b[i];
}).wait();
Step 3 is the whole problem. The compiler has every fact it needs — it can see which queue allocated b, and which queue the kernel was submitted to — and the type system gives it nowhere to record either, so it cannot object. This is the same defect CUDA has with cudaMemcpy taking two void* and a direction enum, and it survives into SYCL because the C++ type system was never asked to carry the device.
The buffer/accessor model, SYCL’s older alternative to USM, avoids this by taking ownership instead. You declare a sycl::buffer, request an accessor inside a command group, and the runtime builds a dependency graph and moves data for you:
1
2
3
4
5
6
7
sycl::buffer<float> buf(host_ptr, sycl::range<1>(n));
q.submit([&](sycl::handler &h) {
// 1. The accessor declares intent (read / write) on a buffer. The runtime
// uses these to order kernels and to insert transfers automatically.
sycl::accessor acc(buf, h, sycl::read_write);
h.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) { acc[i] *= 2.0f; });
});
This is safer, and it is genuinely good design. It also trades one untyped thing for another: the transfers are now correct by construction and completely invisible. Nothing in the source marks where a copy happens, how large it is, or which link it crosses. You find out from a profiler.
So the honest version of the gap is narrower than “SYCL doesn’t type memory,” which is false. It is: SYCL types where memory lives within a device, and does not type which device it lives on, because placement is a runtime property of a queue rather than a static property of a value.
Kokkos is the harder objection
If the SYCL discussion goes well, someone who does HPC for a living asks the better question. Kokkos has had memory spaces as template parameters for a decade.
1
2
3
4
5
6
// The memory space is a template parameter, so these are different types.
Kokkos::View<double*, Kokkos::CudaSpace> d("device", n);
Kokkos::View<double*, Kokkos::HostSpace> h("host", n);
// Explicit transfer. No implicit migration, no hidden copy.
Kokkos::deep_copy(h, d);
That is memory space in the type6, in C++, in production, in codes that run on Frontier and Aurora. And the checking is real. View-to-View assignment across spaces goes through a compile-time trait:
1
2
3
4
template <typename DstMemorySpace, typename SrcMemorySpace>
struct MemorySpaceAccess {
enum { assignable = std::is_same_v<DstMemorySpace, SrcMemorySpace> };
enum { accessible = assignable };
with SpaceAccessibility<AccessSpace, MemorySpace> layered on top to express which execution spaces can reach which memory spaces.7 Write h = d; for incompatible spaces and the program does not compile. RAJA, by contrast, deliberately stays out of memory management and gives you execution policies only, so Kokkos is the one to beat here.
Anyone claiming “nobody puts memory space in the type” has just been refuted by a library that shipped in 2015.
The precise place Kokkos stops, and why
So I went to find where the compile-time guarantee ends, expecting a design compromise. What is actually there is more interesting.
Element access on a View routes through this:
1
2
3
4
5
6
7
8
template <class MemorySpace>
KOKKOS_FUNCTION void runtime_check_memory_access_violation(
SharedAllocationTracker const &track) {
KOKKOS_IF_ON_HOST(((void)RuntimeCheckBasicViewMemoryAccessViolation<
MemorySpace, DefaultHostExecutionSpace>(track);))
KOKKOS_IF_ON_DEVICE(((void)RuntimeCheckBasicViewMemoryAccessViolation<
MemorySpace, DefaultExecutionSpace>(track);))
}
and that helper is selected on a compile-time boolean:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
// 1. Primary template: the memory IS accessible from this execution space.
// Empty constructor, so the check costs nothing at run time.
template <class MemorySpace, class AccessSpace,
bool = SpaceAccessibility<AccessSpace, MemorySpace>::accessible>
struct RuntimeCheckBasicViewMemoryAccessViolation {
KOKKOS_FUNCTION RuntimeCheckBasicViewMemoryAccessViolation(
Kokkos::Impl::SharedAllocationTracker const &) {}
};
// 2. Specialization for accessible == false. This is the illegal case, and
// it is resolved at COMPILE time -- but the action taken is a RUNTIME abort.
template <class MemorySpace, class AccessSpace>
struct RuntimeCheckBasicViewMemoryAccessViolation<MemorySpace, AccessSpace,
false> {
KOKKOS_FUNCTION RuntimeCheckBasicViewMemoryAccessViolation(
Kokkos::Impl::SharedAllocationTracker const &tracker) {
// builds an error message with the View's label, then Kokkos::abort()
}
};
Read what that says. The template is specialized on SpaceAccessibility<AccessSpace, MemorySpace>::accessible, which is a constexpr boolean. At the point the compiler picks the false specialization, it has already proven the access is illegal. And what it emits is a call to Kokkos::abort with a nice error message.
The information is available statically. The diagnostic is deferred to run time anyway.8
That is not Kokkos being careless. It is the only thing a C++ library can do, and the reason is worth understanding because it generalises to every library-based attempt at this. View::operator() is a member function template instantiated from generic code. Whether a particular call executes on the host or inside a device kernel is not part of any function’s type — it is a property of the surrounding lexical context, approximated by macros like KOKKOS_IF_ON_HOST that expand differently in the host and device compilation passes. A static_assert fires when a template is instantiated, not when it is called in a context that would be illegal. Since generic Kokkos code routinely instantiates a View’s accessor in both passes, a static_assert would reject correct programs. The library is forced down to a runtime abort by the language, having done all the reasoning correctly.
The obstacle is that “where this code runs” is not part of a function’s type. Memory space can be a template parameter because it is a property of data. Execution space cannot be attached to a call site the same way, because a C++ function signature has nowhere to put it.
What none of them type
Lay the four systems side by side on the question of what is static:
graph TB
subgraph S1["Typed: memory region within a device"]
A1["OpenCL C: __global / __local / __private"]
A2["SYCL: multi_ptr address_space"]
A3["Kokkos: View<T, CudaSpace>"]
end
subgraph S2["Typed: which physical device"]
B1["nobody"]
end
subgraph S3["Typed: where a function may run"]
C1["nobody -- macros approximate it"]
end
subgraph S4["Typed: cost of moving between them"]
D1["nobody"]
end
style S1 fill:#14532d,color:#fff
style S2 fill:#7f1d1d,color:#fff
style S3 fill:#7f1d1d,color:#fff
style S4 fill:#7f1d1d,color:#fff
The first column is genuinely solved, and has been for years. The other three are open, and they are not open because the problem is hard to notice. They are open because C++ is the substrate, and a library can only type what a template parameter can carry.
Which device: a runtime value everywhere. queue, cl_device_id, cudaSetDevice(int). Two allocations on two GPUs have the same type in all four systems.
Where a function may run: approximated by macros and attributes — __device__, KOKKOS_FUNCTION, KOKKOS_IF_ON_HOST. These are preprocessor conditionals and calling-convention annotations, not types. You cannot write a signature that requires its caller to be executing on a particular device, which is exactly why Kokkos ends at a runtime abort.
What a transfer costs: nothing models it. deep_copy between two spaces is one call whether it crosses a PCIe link, NVLink, or a network. The buffer/accessor model will insert transfers for you without telling you what they cost. A cost model needs a topology, and none of these type systems has a notion of topology at all.
What I am not claiming
Two guardrails, because the failure mode of a post like this is overclaiming and the audience for it will notice.
Nothing above says SYCL or Kokkos is badly designed. Both hit their actual goals. SYCL set out to make heterogeneous programming possible in standard C++ across vendors, and it did. Kokkos set out to give HPC codes performance portability without a rewrite per machine, and it did, on the largest machines that exist. Judged against what they were for, the criticism does not land.
And nothing above says the missing pieces are free. Putting placement in the type system means a function’s signature grows a component that every caller must satisfy, which means inference, which means the compiler needs a topology model, which means the topology becomes part of your build configuration. That is a real cost and possibly a bad trade for a codebase that runs on one machine. Kokkos users who target one supercomputer at a time are not obviously wrong to prefer the runtime abort.
The claim is narrower. The reason placement is untyped everywhere is not that it was considered and rejected. It is that every one of these systems is a library, and a library cannot introduce a static property that the host language has no place to record. Kokkos proves the point precisely because it went as far as C++ permits and then stopped at a Kokkos::abort it had already proven unnecessary.
That is a constraint of the substrate, not of the idea. What a language could do with it — and whether the cost is worth paying — is a separate argument, and one I will make separately.
References
Disclaimer: Researched and drafted with AI assistance (Claude Opus 5). Direction, technical judgment, and final edits are mine; every claim is traceable to the sources cited above. The SYCL and Kokkos declarations quoted here were read from the DPC++ and Kokkos sources at the links given; I have not executed the example code in this post.
AdaptiveCpp. An independent SYCL implementation (formerly hipSYCL / Open SYCL) targeting CUDA, HIP, Level Zero and host OpenMP, which is the practical evidence that SYCL is not a single-vendor interface. (Project) ↩︎
SYCL
access::address_space. The enum as declared in DPC++, including the SYCL 2020 deprecation ofconstant_spaceand the addition ofgeneric_space. (sycl/include/sycl/access/access.hpp) ↩︎sycl_ext_intel_usm_address_spaces. The Intel extension addingdeviceandhostas subsets of the global address space, withdevice_ptrandhost_ptrinterfaces onmulti_ptr. The document states the goal is to let users “explicitly tell the compiler which address space a pointer resides in for optimization purposes,” and that using these objects “allows the compiler to perform better alias analysis, which typically leads to better throughput and smaller silicon area.” (intel/llvm) ↩︎ ↩︎2SYCL device selectors.
default_selector_v,gpu_selector_v,cpu_selector_vandaccelerator_selector_vare declared as functions taking aconst device &and returning anintscore; the SYCL 1.2.1device_selectorbase class is deprecated in SYCL 2020. (sycl/include/sycl/device_selector.hpp) ↩︎SYCL USM allocation.
malloc_device,malloc_hostandmalloc_shared, in bothvoid*and templatedT*forms, taking aqueueor adevice+context. No overload returns a type parameterised by the device. (sycl/include/sycl/usm.hpp) ↩︎Kokkos
View. The class template and its memory-space template parameter; when omitted,DefaultExecutionSpace::memory_spaceis used. (Kokkos wiki) ↩︎Kokkos
SpaceAccessibilityandMemorySpaceAccess. The compile-time traits expressing which execution spaces may access which memory spaces, and which View assignments are permitted. (core/src/Kokkos_Concepts.hpp) ↩︎Kokkos View access checking.
runtime_check_memory_access_violationandRuntimeCheckBasicViewMemoryAccessViolation, whosefalsespecialization is selected on theconstexprSpaceAccessibility<...>::accessibleand emits aKokkos::abort. (core/src/View/Kokkos_ViewAccessPreconditionsCheck.hpp) ↩︎