What happens
On dgx (GB10, sm_121), dispatching any vt:: op that has a CPU kernel but
no native CUDA kernel segfaults instead of either running or throwing.
Reproduced 2026-08-13 while landing the RED half of W2 of
.agents/specs/mamba2-ssd.md (#496): the three
Mamba2 SSD ops had CPU kernels and, at that commit, no CUDA kernels. Every CUDA
test case died at the first dispatch:
[vt reference-tier] op=113 device=1 has NO native kernel; running the PORTABLE CPU fallback (correct but slow)
...
FATAL ERROR: test case CRASHED: SIGSEGV - Segmentation violation signal
[doctest] Status: FAILURE!
Same for op=114 and op=115 (kMamba2StateUpdate, kRmsNormGatedGroup); exit
code 139 on all three binaries.
Why
ReferenceTierEligible (src/vt/op_provider.cpp:515-526) gates the portable CPU
reference tier on Backend::UnifiedMemory():
Backend* b = TryGetBackend(device);
return b != nullptr && b->UnifiedMemory();
and its contract in include/vt/op_provider.h states the invariant as "correct
ONLY where host and device memory alias (Metal StorageModeShared, GB10 /
integrated Vulkan, CPU) — Backend::UnifiedMemory()".
CudaBackend reports UnifiedMemory() == true on GB10
(caps.pageable_memory_access && caps.integrated, src/vt/cuda/cuda_backend.cu
Registrar), so the tier installs. But CudaBackend::Alloc is a plain
cudaMalloc (cuda_backend.cu:77-79), and a cudaMalloc pointer is not host
dereferenceable on GB10 — which include/vt/backend.h already says in as many
words, on the doc comment of the other predicate:
CUDA on GB10 reports unified memory because host and device address the same
physical RAM, yet a plain cudaMalloc pointer is still not
host-dereferenceable.
So the two contracts contradict each other, and the segfault settles which one is
right. Backend::DeviceMemoryIsHostAddressable() — which CudaBackend does not
override, so it is false — is the predicate that actually describes the
requirement ("a pointer returned by Alloc() may be DEREFERENCED BY THE HOST
directly"). ReferenceTierEligible is testing the weaker one.
Why this is worse than a crash
The tier is meant to be the safety net that lets a partial backend run
(forward_native's equivalent). On the CUDA gate box it is instead a latent
segfault on every missing kernel, and — because it announces itself as
"correct but slow" — it reads as a working fallback right up to the fault. Any
new CUDA op lands with this trap under it.
It also nearly produced a false green: a device test whose op silently ran on
the host CPU would have passed every numeric assertion while nothing ran on the
GPU. The W2 suites now assert the selected provider is native
(GetOpProviderStats(op, kCUDA).last_selected != kReferenceProviderName), which
is how the fallback was noticed at all.
Suggested fix (needs its own row/spec — NOT done here)
Gate on DeviceMemoryIsHostAddressable() rather than UnifiedMemory(), so CUDA
answers false and a missing kernel throws the way it did before the tier
existed. Metal (StorageModeShared) and integrated Vulkan (persistently mapped
HOST_VISIBLE) both already opt into the narrower predicate, so they keep the
tier.
That changes the semantics of a shared seam three backends depend on, so it
takes the normal row / spec / fresh-review path rather than an in-flow fix, and
is deliberately not part of #496's W2. Recorded here so it is not
rediscovered by the next CUDA op that lands without a kernel.
Evidence
- Box:
promaxgb10-4ad8, NVIDIA GB10, CUDA 13.0.88, Release build,
-DVLLM_CPP_CUDA_ARCHITECTURES=121a, CUTLASS 4.5.0, FA2 enabled.
- Binaries:
test_ops_mamba2_ssd, test_ops_mamba2_state_update,
test_ops_mamba2_gated_norm — all exit 139 at the first CUDA dispatch.
- Anchors:
src/vt/op_provider.cpp:515-526, include/vt/op_provider.h
(reference-tier section), include/vt/backend.h
(DeviceMemoryIsHostAddressable), src/vt/cuda/cuda_backend.cu:77-79 and its
Registrar.
What happens
On
dgx(GB10, sm_121), dispatching anyvt::op that has a CPU kernel butno native CUDA kernel segfaults instead of either running or throwing.
Reproduced 2026-08-13 while landing the RED half of W2 of
.agents/specs/mamba2-ssd.md(#496): the threeMamba2 SSD ops had CPU kernels and, at that commit, no CUDA kernels. Every CUDA
test case died at the first dispatch:
Same for op=114 and op=115 (
kMamba2StateUpdate,kRmsNormGatedGroup); exitcode 139 on all three binaries.
Why
ReferenceTierEligible(src/vt/op_provider.cpp:515-526) gates the portable CPUreference tier on
Backend::UnifiedMemory():and its contract in
include/vt/op_provider.hstates the invariant as "correctONLY where host and device memory alias (Metal StorageModeShared, GB10 /
integrated Vulkan, CPU) —
Backend::UnifiedMemory()".CudaBackendreportsUnifiedMemory() == trueon GB10(
caps.pageable_memory_access && caps.integrated,src/vt/cuda/cuda_backend.cuRegistrar), so the tier installs. But
CudaBackend::Allocis a plaincudaMalloc(cuda_backend.cu:77-79), and acudaMallocpointer is not hostdereferenceable on GB10 — which
include/vt/backend.halready says in as manywords, on the doc comment of the other predicate:
So the two contracts contradict each other, and the segfault settles which one is
right.
Backend::DeviceMemoryIsHostAddressable()— whichCudaBackenddoes notoverride, so it is
false— is the predicate that actually describes therequirement ("a pointer returned by Alloc() may be DEREFERENCED BY THE HOST
directly").
ReferenceTierEligibleis testing the weaker one.Why this is worse than a crash
The tier is meant to be the safety net that lets a partial backend run
(
forward_native's equivalent). On the CUDA gate box it is instead a latentsegfault on every missing kernel, and — because it announces itself as
"correct but slow" — it reads as a working fallback right up to the fault. Any
new CUDA op lands with this trap under it.
It also nearly produced a false green: a device test whose op silently ran on
the host CPU would have passed every numeric assertion while nothing ran on the
GPU. The W2 suites now assert the selected provider is native
(
GetOpProviderStats(op, kCUDA).last_selected != kReferenceProviderName), whichis how the fallback was noticed at all.
Suggested fix (needs its own row/spec — NOT done here)
Gate on
DeviceMemoryIsHostAddressable()rather thanUnifiedMemory(), so CUDAanswers false and a missing kernel throws the way it did before the tier
existed. Metal (StorageModeShared) and integrated Vulkan (persistently mapped
HOST_VISIBLE) both already opt into the narrower predicate, so they keep the
tier.
That changes the semantics of a shared seam three backends depend on, so it
takes the normal row / spec / fresh-review path rather than an in-flow fix, and
is deliberately not part of #496's W2. Recorded here so it is not
rediscovered by the next CUDA op that lands without a kernel.
Evidence
promaxgb10-4ad8, NVIDIA GB10, CUDA 13.0.88, Release build,-DVLLM_CPP_CUDA_ARCHITECTURES=121a, CUTLASS 4.5.0, FA2 enabled.test_ops_mamba2_ssd,test_ops_mamba2_state_update,test_ops_mamba2_gated_norm— all exit 139 at the first CUDA dispatch.src/vt/op_provider.cpp:515-526,include/vt/op_provider.h(reference-tier section),
include/vt/backend.h(
DeviceMemoryIsHostAddressable),src/vt/cuda/cuda_backend.cu:77-79and itsRegistrar.