docs / specification / xpu

Cajeta XPU capability matrix — NVIDIA · AMD · Vulkan · Core

A per-feature matrix of the Cajeta GPU-compute (XPU) backend across the three target platforms, plus a Core column for what the shared layer provides regardless of vendor. The point of the matrix is to make the variance surface legible: where all three have a native primitive, where one lacks it but Cajeta supplies an abstraction, and where an abstraction is not cleanly possible (with the reason).

This is a companion to the backend-variance discipline in CajetaXPU-Variance.md; it extends the original NVIDIA∩AMD two-backend reckoning with the Vulkan/SPIR-V column. The stable launch + kernel-arg FFI contract these backends dispatch through (the register trio, the argv marshalling per parameter kind, per-launch device targeting, and the ABI version policy) is frozen in CajetaXPU-FFI.md.

A fourth backend — CPU — and a runtime dispatcher now exist (see CajetaCPU.md). The CPU is deliberately not given a column here: the three columns below measure the GPU variance surface, and the CPU is a different shape — it has no hardware grid or coordinate intrinsics, so its sole seam fork is the coordinate source (the kernel gains 9 trailing i32 coordinate params; KernelThread/Workgroup reads pull from those — the grid→threads model), buffers are flat addrspace(0), the wave is the host SIMD vector (Inc 5C), and workgroup barriers run via work-item loop fission (Inc 6 — split the kernel at each barrier, loop each region over the block; per-block shared memory + context arrays; Inc 7 — multiple barriers per uniform loop; Inc 8 — nested uniform loops with barriers; Inc 9 — wave ops + barriers composed; Inc 10 — dynamic-sized shared memory). The body walk is the same ~90% Core. Separately, the runtime dispatcher (in the C runtime, the sole launch path for compiled programs) selects among bundled+available backends at startup (CUDA → HIP → Vulkan → CPU) and routes launch + KernelBuffer<T> device memory to the winner — so “run anywhere, degrade to CPU” is ordinary dispatch reaching its guaranteed terminal. The one launch-ABI asymmetry it exposed is Vulkan’s: the uniform kernelParams argv (buffer handles + raw scalars) doesn’t map to Vulkan’s descriptor-set model, so scalars are wrapped in transient SSBOs from per-kernel parameter metadata (rows in §2 already note the descriptor-set fork).

Provenance & how to read this

The three platform columns do not have equal evidentiary weight yet, and the matrix says so explicitly rather than pretending otherwise:

Build-discovered correction (2026-05-30). An earlier draft of this matrix projected the Vulkan column on the assumption that KernelBuffer Device Address would carry buffers as raw pointers (keeping the kernel body identical to NV/AMD). The build overturned that: LLVM 23’s SPIR-V backend has no PhysicalStorageBuffer/BDA path from IR, so the only Vulkan buffer model is descriptor-set storage buffers, which fork the kernel signature and the body’s buffer access — a bigger fork than AMD. The cells below reflect what was actually built and measured, not the original projection.

Cell vocabulary

• native — the platform has a direct hardware/IR primitive; no abstraction layer needed.
• abstraction: … — no direct primitive, but Cajeta can synthesize the capability cleanly; the … says how.
• forks: … — supported, but the implementation diverges from the others in a way that lands on the seam; the … says where.
• not possible: … — cannot be provided on this platform as specified; the … says why.
• — — not applicable / not yet a capability on any platform.

⚑ marks a Vulkan cell that carries a measured caveat or limitation the build surfaced (the fixed compile-time block dim, deferred dynamic shared) — see the status section. (The barrier limitation the build found is now fixed by a post-emit pass — §1.)

The Core column describes what the shared KernelLowering walk + the backend-neutral registration provide independent of vendor — the ≈90% that did not fork between NVIDIA and AMD, and which the Vulkan analysis expects to hold for a third time.

1. Execution model — coordinate reads & barriers

These are the entire LoweringTarget leaf-read interface (src/cajeta/xpu/lowering/LoweringTarget.h). The kernel-body AST walk that consumes them is shared and does not fork.

Reading: every coordinate read is native on Vulkan, and two of them (workgroupDim, globalId) are native single-intrinsic reads where AMD needed a structural workaround — on this axis Vulkan is the least divergent of the three. The one wrinkle is the barrier: LLVM 23 emits Vulkan-forbidden semantics, so a one-instruction post-emit fixup (see the cell above) rewrites it — after which the execution-model column is fully native + spec-valid on all three backends.

2. Address spaces & memory model

Numeric address spaces from src/cajeta/xpu/core/AddressSpace.h. NVIDIA and AMD share the LLVM numbering; SPIR-V uses storage classes (the integers below are the Khronos storage-class enum values, not LLVM address spaces).

Reading: address space is a near-non-fork across all three (NV/AMD identical; Vulkan’s only real difference is the absence of a generic flat space, which the explicit-AS lowering sidesteps). The single alloca fork (0/5/7) stays a one-line seam method.

3. Kernel-argument ABI — the Vulkan fork (signature + buffer access)

NVIDIA and AMD both pass kernel args as a flat kernelParams array: buffers as raw addrspace(1) device pointers, scalars by value. Vulkan has no raw-pointer kernel-argument ABI. This is the architectural fork the Vulkan bring-up turns on.

Reading: Vulkan’s missing raw-pointer ABI is the headline divergence — and the build proved it’s not bridgeable by KernelBuffer Device Address (no IR path), only by descriptor-set SSBOs. That is the “bigger fork than AMD” the pass set out to measure: where AMD forked only the coordinate leaf reads, Vulkan forks the kernel signature (void main(), descriptor binds) and the body’s buffer element access (getpointer vs GEP). The coordinate reads, control flow, operators, scalar-slot model, and addrspace(3) shared memory still stay shared — so the core is smaller than NV∩AMD’s ≈90%, but still the majority.

4. Launch / dispatch / runtime

Reading: the launch-config shape (grid/block + a buffer/scalar arg list) stays universal at the Cajeta surface, but block dim partially migrates from launch-time (NV/AMD) to pipeline-creation-time (Vulkan), and the runtime driver forks heavily. The fork is contained to the driver layer; the language-level kernel.launch(...) syntax stays backend-neutral.

5. Codegen pipeline (triple · assembler · binary format)

Reading: Vulkan is the only one of the three that needs no external assembler or linker — the TargetMachine produces the shippable binary directly.

6. Registration & frontend (the backend-neutral core)

Reading: the entire frontend + the name-keyed registration are vendor-neutral and were never touched when AMD landed; the Vulkan analysis expects the same — the third backend inherits all of it for free.

7. Shared (workgroup) memory

Reading: static shared is a clean native three-way. Dynamic shared is the second place (after block dim, §4) where a value Cajeta treats as launch-time on NV/AMD becomes pipeline-creation-time on Vulkan — bridgeable by a spec constant + pipeline cache, but it cannot be a free per-dispatch scalar, and the matrix says why.

8. Wave / subgroup ops (the @Wave feature — readlane + ballot + reduce built & measured)

Wave ops are the archetypal variance-shaped feature and the headline test of the seam — and the seam held. Wave.shuffleSync (readlane / shuffle-by-index), Wave.ballotSync, Wave.reduceSum, and Wave.laneId lower through the shared AST walk with only five LoweringTarget methods (waveWidth/waveShuffle/waveBallot/ waveReduceSum/waveLaneId) forking. Built 2026-05-30; emit-verified on all three GPU backends, run on-device on AMD + Vulkan + NVIDIA (RTX 4090, 2026-06-16 — shuffle/ballot/reduceSum/reduce-family/laneId/width/rotate/prefix-scan); wave = SIMD lane on the CPU backend (Inc 5C, 2026-05-31).

Reading: wave ops are native on all three. The reduce probe overturned its own hypothesis. The guess (recorded here in the prior pass) was that reduce would invert comprehensiveness — one native intrinsic on Vulkan vs. a shuffle/DPP butterfly sequence on NV/AMD. The build showed the opposite: all three expose a single hardware wave-reduce intrinsic in LLVM 23, so reduce maps as cleanly as shuffle/ballot (waveReduceSum is a one-liner per backend). The real asymmetries are narrower and were only visible by running it:

• NVPTX redux.sync is gated on sm_80+ (Ampere); below that a butterfly-shuffle fallback would be needed (out of scope — the emit test targets sm_89).
• AMDGPU folds a uniform-constant operand to wave.reduce.add back to the operand instead of summing it (reduceSum(1) returned 1 on-device); a divergent operand takes the real reduction path. The device test feeds a buffer of 1s so the reduction actually runs.
• Ballot shape still forks (NV i32, AMD i32 wave32, Vulkan <4 x i32> 128-bit — the Core API normalizes to i64).

Tests: XpuWaveEmitTests (3 backends + spirv-val, shuffle/ballot/reduce), XpuWaveDeviceTests (shuffle/ballot/reduce on-device on AMD + Vulkan + NVIDIA; the nvptx* arms cover shuffle/ballot, reduceSum, laneId, width, rotate, the reduce family, and prefix scans; the reduce/width checks are width-agnostic — sum of 1s over a full wave == wave width ∈ {32, 64}, which is 32 on NVIDIA).

9. Atomics & memory ordering

10. Deferred / not-yet capabilities (no backend has these)

Tracked here so the matrix is honest about the frontier, not just the floor. These are backend-neutral gaps, unaffected by the Vulkan column.

11. The abstraction-layer ledger (the user’s core question)

Where a platform lacks a native primitive — can Cajeta provide it, and if not, why?

Provided cleanly (abstraction layer bridges the gap):

• Vulkan · raw-pointer buffer args → descriptor-set SSBOs (resource.handlefrombinding + getpointer); scalars as single-element SSBOs. Forks signature + body buffer access; the rest of the body stays shared. (§3)

• Vulkan · workgroup barrier → group.memory.barrier.with.group.sync + a one-instruction post-emit fixup (SpirvBackend::fixupControlBarriers) that corrects LLVM 23’s Vulkan-forbidden SequentiallyConsistent semantics to WorkgroupMemory|AcquireRelease. Now passes strict spirv-val and runs on-device. (§1)

• Vulkan · block dim → fixed compile-time LocalSize (first cut); spec-constant LocalSizeId is the refinement. (§4)

• AMD · workgroup dim read → dispatch-packet load (no ntid intrinsic, but the value is recoverable). (§1)

• All · tex.sample(sampler, u, v) (Item 8) → one LoweringTarget::sampleTexture seam, four native realizations: CPU C bilinear (__cajeta_xpu_cpu_tex_sample); Vulkan llvm.spv.resource.samplelevel → OpImageSampleExplicitLod (image + sampler descriptors); AMD __ockl_image_sample_2D → image_sample (ROCm device-lib hybrid-linked only for sampling kernels — reuses ROCm’s SRD build + coord normalization rather than hand-packing gfx descriptors); NVIDIA llvm.nvvm.tex.unified.2d → PTX tex.2d. Texture marshals like a KernelBuffer handle, Sampler like a by-value POD; on AMD/NVIDIA the sampler state rides the texture object (built per-launch). CPU+Vulkan+AMD on-device, NVIDIA emit-only. (needs LLVM 23 for Vulkan)

• Texture dimensions — Texture1D/2D/3D/2DArray + fetch/sample work on CPU + Vulkan/RADV + AMD/gfx1151 on-device (NVIDIA emit-only). TextureCube sample and mipmapped Texture2D + explicit-LOD (sampleLod/fetchLod, incl. trilinear cross-level blend) now work on CPU + Vulkan/RADV + AMD/gfx1151 on-device (NVIDIA emit-only). HIP’s high-level array APIs don’t support these on gfx1151 — hipMalloc3DArray[hipArrayCubemap] → invalid-arg, hipMallocMipmappedArray → hipErrorNotSupported(801), unchanged across ROCm 7.11.0 (a documented HIP gap — the official “Texture Management [Not supported]” group — not a hardware limit). cajeta emulates both on AMD: cubemaps via a LAYERED hipArray + in-kernel face projection, and mipmaps via a hand-built gfx11 image SRD over an addrlib-tiled hipMalloc (option B), sampled through the unchanged __ockl_image_sample_lod_2D seam — proven bit-exact on-device (plans/gpu/xpu/probes/mipprobe.cpp). The mip emulation rides the optional libcajeta_amdtex helper (vendored addrlib, dlopen’d like libamdhip64); where it’s absent or the gfx arch isn’t in its config table the AMD mip path degrades to unsupported and the Vulkan backend remains the fallback. Scope: explicit-LOD only (no auto-LOD/derivatives — N/A in compute), no anisotropic, no seamless cube edges; v1 mip format = R32F. (XpuHipDispatchDeviceTests.{mipmapFetchAndSampleLod,mipTrilinearBlend,textureCubeSample}RoutesToHipOnDevice; project memory project_amd_mip_emulation_shipped)

Not cleanly possible (and why):

• Vulkan · KernelBuffer Device Address (raw pointers) → LLVM 23’s SPIR-V backend has no PhysicalStorageBuffer/BDA intrinsic from IR. This forced the descriptor-set model above — the central build-discovered finding. (§3)
• Vulkan · generic/flat address space → SPIR-V has no Generic storage class outside the OpenCL Kernel capability. Benign: the lowerer tracks every pointer’s address space explicitly and never relies on flat generic deref. (§2)
• Vulkan · dynamic shared as a free per-dispatch byte count → workgroup arrays are sized at pipeline creation, not per-vkCmdDispatch. Deferred this pass. (§7)

Needed nowhere (already shared across all three): the entire kernel-body AST walk, the coordinate leaf reads, addrspace(3) shared memory (→ Workgroup on SPIR-V, a non-fork), the name-keyed registration symbol, and the whole frontend.

Status — Vulkan column now measured

The Vulkan backend was built 2026-05-30 by threading SPIR-V through the same LoweringTarget seam AMD uses. All increments landed; the on-device tests run on the Strix Halo APU via RADV.

One cell carries a measured caveat — block dim is fixed at SPIR-V compile time (§4); dynamic shared is deferred. The barrier limitation found during the build is now fixed (§1). Everything else in the Vulkan column is on-device-measured and spec-valid.

Part C cutting-edge SPIR-V (via the cajeta-llvm fork, LLVM 23)

Both delivered on-device on the RADV / STRIX_HALO (Radeon 8060S) box (2026-06-04), through the fork’s cajeta-spirv branch. The native matrix-core path is now wired on all three GPU backends (Vulkan OpCooperativeMatrix, AMD RDNA3 WMMA, NVIDIA tensor-core wmma, 2026-06-16); the native RT-core ray-query seam is on-device- validated on RADV (Linux) + the RTX 4090 (Windows) via Vulkan, and on NVIDIA CUDA via OptiX (2026-06-17) — only the AMDGPU backend stays software-BVH-only. NVIDIA now runs both via the portable/software tiers on-device (RTX 4090, 2026-06-16): cooperative-matrix on the portable flat-tile tier (f16/f16→f32 bit-exact), and ray-query over the portable software BVH (the same SoftwareRayQuery walk the CPU uses — AABB/triangle/nearest/barycentrics/front-face all match the CPU results). NVIDIA cooperative-matrix now ALSO runs on the native tensor-core (wmma) tier — 16×16×16 f16→f32 row-major, bit-exact on the RTX 4090 via the NVVM wmma.load/wmma.mma/wmma.store intrinsics (the cajeta seam derives the fragment struct type from the load intrinsic so load/mma/store agree; the warp- collective launch uses a full warp). So all three GPU backends now have a native cooperative-matrix path (Vulkan OpCooperativeMatrix, AMD RDNA3 WMMA, NVIDIA tensor-core wmma). The native RT-core ray-query path is now on-device-validated on the RTX 4090’s Windows Vulkan backend (2026-06-16): AABB + triangle BLAS/TLAS traced via OpRayQuery on the RT cores, results matching the software-BVH oracle, with AUTO resolving to native and AsImpl.Software/AsImpl.Native selectable. NVIDIA CUDA also reaches its RT cores via OptiX (2026-06-17): a RayQuery @Kernel against an OptiX-impl AS lowers to a separate OptiX program-PTX module (NvptxOptixRayQuery) dispatched by optixLaunch — four canonical shapes (AABB candidate-count, triangle nearest-hit, triangle candidate getters, committed-triangle / front-face) match the software oracle (777) on the 4090. Under M3, AUTO on CUDA prefers the RT cores (Device.supports(Capability.RayQueryRtCore)): the AS is a multi-impl noun that retains the software floor and builds the OptiX rep lazily on the first supported-shape launch, with launch-time selection routing Unsupported shapes back to the floor (no fault, no opt-in — see docs/gpu/RayQuery.md §5–§6). The software BVH stays the portable floor (CPU/AMD, non-RT Vulkan, and any Unsupported-shape kernel). AMD gets the symmetric ray-query software-BVH path (AmdgpuTarget.accelImpl() == SoftwareBvh

• a HIP noun provider) — code-complete + compile-verified, on-device-PENDING the gfx1151 box (both NVPTX and AMDGPU had the same latent accelImpl() gap that made RayQuery throw; now fixed). AMD’s cooperative-matrix is already native WMMA.

float16 is IEEE binary16 (getHalfTy, not bfloat) to match the device’s VK_COMPONENT_TYPE_FLOAT16_KHR cooperative-matrix config.

Ray-query AS impl, by platform. The AccelerationStructure noun has two impls (CajetaGPU §1.5): the native Vulkan BLAS+TLAS (OpRayQuery, RT cores) and the portable software BVH (the SoftwareRayQuery walk over a KernelBuffer<float32>, shared with CPU/NVPTX/AMD). On any ray-query-capable Vulkan device — RADV on Linux AND the RTX 4090 on Windows — AUTO resolves to native; the caj_native_ rayquery_available resolver and the RayQueryNative capability share one condition (active backend is Vulkan + the device advertises ray query), so they never disagree. On a non-RT Vulkan device (or the CPU/NVPTX/AMD backends, which have no native inline ray-query seam) AUTO resolves to the software BVH floor. The impl is selectable per AS: AccelerationStructure.of(.., AsImpl.Native | AsImpl.Software) or the process-wide CAJETA_GPU_AS_IMPL=native|software override; the chosen impl is recorded on the noun (AccelerationStructure.implTag()) and the verb follows it.

CUDA gets a third impl — OptiX (implTag 2, 2026-06-17). NVIDIA has no inline ray-query seam, so its RT cores are reached through an OptiX pipeline that the compiler emits as a lowering of the same inline RayQuery verb (NvptxOptixRayQuery). M3 made AUTO prefer OptiX safely via a multi-impl noun + launch-time selection: an AccelerationStructure always retains the software-BVH floor and (on CUDA, under AUTO) builds the OptiX rep lazily on the first supported-shape launch — implSet() reports the live set (1<<impl per rep; e.g. Software|OptiX = 5), while implTag() keeps the single-primary read. At launch the runtime picks, per consuming kernel, the best impl it can traverse (a registered OptiX program + the AS carrying OptiX → optixLaunch; else the software cubin over the floor — an Unsupported shape can no longer be handed an OptixAs* to misread). =optix forces eager OptiX (+ retained floor); =software forces the floor. Device.supports(Capability.RayQueryRtCore) reports the OptiX path per device (distinct from RayQueryNative, false on CUDA). This supersedes the earlier opt-in-only / deferred- flip policy (the fault that made it unsafe is gone). See documents/gpu-rayquery-optix/ rayquery-optix-m3-multiimpl-{spec,plan}.md.

On-device (RTX 4090, Windows, 2026-06-16): native AABB + triangle ray query — hits/primitive-index/T/barycentrics/front-face/nearest-hit — all match the software oracle; autoRecordsNativeImplOnDevice proves AUTO records native (impl 1) on the 4090, and forcedNativeOfApiOnDevice/forcedSoftwareOfApiOnDevice prove both impls are selectable and agree. Bugs fixed on the way to native: the native triangle builder recorded a BLAS as the traceable AS (must be a TLAS — NVIDIA returns no hits, RADV tolerated it; now wraps the BLAS in a TLAS like the AABB path), and the runtime resolver + RayQueryNative capability were both Win32-gated (now removed — the 4090’s native RT-core path is wired and validated). AMD’s HIP backend still uses the software BVH. NVIDIA’s CUDA backend is growing an OptiX RT-core tier (CAJ_AS_IMPL_OPTIX): NVIDIA RT cores are reached via OptiX (a pipeline model, no inline-RQ intrinsic), so M0 proved OptiX results match the software oracle on the 4090 and M1 landed the runtime AS provider (build + record the OptiX AS on-device, implTag 2); the traversal verb (NVPTX→OptiX codegen + optixLaunch) is in progress (M2). OptiX is a compile-time-only dependency — the engine is the driver’s nvoptix.dll. See documents/gpu-rayquery-optix/.

Generated 2026-05-30; updated 2026-06-04 with Part C cooperative-matrix (CM1–CM5 + GEMM) and ray-query (Inc 1–3c) on-device landings. Updated 2026-06-16: the NVIDIA compute substrate is now on-device-measured on an RTX 4090 (sm_89, native Windows CUDA — the B5 WSL2 runner turned out to be unnecessary). Two real bugs were found and fixed by the on-device bring-up: (1) NvptxRegistration never emitted the __cajeta_xpu_register_kernel_params ctor (AMD/Vulkan did), so find_kparams returned NULL on CUDA and bindless/texture/image arg translation silently no-op’d; (2) __cajeta_xpu_register_module dropped re-registration of an existing kernel name (first-wins) instead of overwriting + resetting the cached module — a use-after-free for a second JIT’d program reusing a kernel name. Both fixed; all three GPU backends + CPU are now on-device-measured for the compute substrate. Follow-up landed the same day: NVIDIA wave ops, the texture+surface runtime, ray-query (portable software BVH — new NvptxTarget.accelImpl() == SoftwareBvh + a CUDA noun provider), and cooperative-matrix (portable flat-tile tier) are now all on-device-verified on the RTX 4090. What remains future on NVPTX is only the native tensor-core (wmma/mma.sync) and RT-core seams — the portable/software tiers run today and match the CPU/Vulkan results. See CajetaXPU-Variance.md for the NVIDIA∩AMD variance discipline and cajeta-gpu-plan.md Part C.

Source: docs/specification/xpu/CajetaXPU-Matrix.md · 29 min read · 6354 words