Half/bf16 CUDA atomics will need CAS fallbacks — and a 4-byte alignment precondition on user DLPack pointers

[balbasty]

Found by: claude-jitfields-to-fastfields, during the atomics audit (#97). Filed before any implementation exists, because this is the kind of precondition that otherwise gets discovered by a segfault months later.

If half/bf16 support lands (#92), the CUDA scatter paths will need atomicAdd on 16-bit types. Native support requires sm_70+ for __half and sm_80+ for __nv_bfloat16, and the target matrix reaches below both:

build	floor	half	bf16
cu11	Kepler/Maxwell	needs CAS fallback	needs CAS fallback
cu12	Volta (sm_70)	native	needs CAS fallback
cu13	Turing (sm_75)	native	needs CAS fallback

So all three tiers need a bf16 fallback and cu11 needs one for half. Today anyAtomicAdd(half*, half) is a hard compile error — gpuAtomicAdd has float and double overloads only.

The machinery to build on exists: AtomicAddIntegerImpl<T,2> in atomic.h, and the shape to copy is upstream PyTorch's AtomicFPOp<at::Half> / AtomicFPOp<at::BFloat16>, guarded on __CUDA_ARCH__ < 700 and < 800 respectively.

The precondition that must be written down first

A 2-byte CAS reads and writes the enclosing aligned 4-byte word. There is no 2-byte atomicCAS on CUDA; the standard technique loads the containing 32-bit word, splices the 16-bit lane, and CAS-es the whole word back.

PyTorch gets away with this because its caching allocator over-aligns every tensor. fastfields takes user-supplied DLPack pointers. For a 16-bit tensor with an odd number of elements whose final element ends the allocation, splicing the enclosing word is a 2-byte out-of-bounds access — a read and a write past the end of memory the caller owns.

Two acceptable resolutions, and the choice should be deliberate:

1. Require 4-byte alignment of data + byte_offset for 16-bit dtypes at the dispatch boundary, and reject otherwise with a clear error. Cheap, but rejects legitimate views.
2. Handle the tail element separately — a non-atomic path or a masked CAS for the final unaligned lane.

Why this cannot be caught by existing validation

• ASan on the CPU will never see it: this is CUDA device code.
• Compile+link cannot see it: it is a runtime address computation.
• The CPU suite cannot see it: there is no CPU counterpart to inherit correctness from.

Only a GPU run under compute-sanitizer --tool memcheck would catch it, and there is no GPU in CI. That is the same structural blind spot as #86 (unzeroed accumulation buffer, wrong sign) and #88 (stream ordering) — all silent-wrong-answer or out-of-bounds classes with no compile-time signature.

Related

• prototype: port teeny's portable half/bfloat16 to C++11 (core/half.h) #92 — the half/bf16 prototype (draft). The header ports cleanly to C++11; this is a constraint on the atomics work that would follow, not on the header.
• atomic.h: integer CAS specialisations retyped from upstream — signed char corrupts neighbouring bytes, unsigned long truncates on Win64 #98 — the existing integer CAS specialisations, one of which (<T,1>) already has a byte-splicing bug of exactly this family.
• CUDA launchers rely on cudaFree's implicit device-wide sync for stream ordering #88, cuda(mesh): add the sdt_naive host launcher; zero the vertex normals #86 — the same validation blind spot.