atomic.h: integer CAS specialisations retyped from upstream — signed `char` corrupts neighbouring bytes, `unsigned long` truncates on Win64

[balbasty]

Found by: claude-jitfields-to-fastfields, during the atomics audit (#97). Reported rather than fixed, because both are currently unreachable and the fix should land with whoever makes them reachable.

include/fastfields/impl/kernels/atomic.h's AtomicAddIntegerImpl<T,N> is derived from PyTorch's THCAtomics.cuh / ATen/cuda/Atomic.cuh, but two of the four size specialisations were retyped during the copy. Upstream has used the fixed-width spellings since at least v1.5.

specialisation	upstream	this copy	consequence
<T,1>	static_cast<uint8_t>	static_cast<char>	char is signed on x86. A value ≥ 0x80 becomes negative, sign-extends when widened into unsigned int newval, and (newval << shift) then clobbers the neighbouring bytes of the enclosing 4-byte word.
<T,8>	static_cast<uint64_t>	static_cast<unsigned long>	unsigned long is 32-bit under LLP64 (Win64) → silent truncation of the upper half.

<T,2> (unsigned short) and <T,4> (unsigned int) are equivalent to upstream and fine.

Currently unreachable — which is the danger

GPU_ATOMIC_INTEGER is never invoked, and every integer gpuAtomicAdd overload in the file is commented out. So neither defect can fire today.

They are traps armed for whoever uncomments them — and adding integer dtype support is exactly the change that would do so. The <T,1> bug in particular corrupts adjacent elements rather than the target, so it would present as data corruption in a neighbouring voxel, not as a wrong result where you were looking.

Fix

Use the fixed-width types, matching upstream:

static_cast<uint8_t>(...) // <T,1>
static_cast<uint64_t>(...) // <T,8>

Related

• cuda_switch.h: NVRTC int64_t/uint64_t defined as long, which is 32-bit on Windows (LLP64) #93 — cuda_switch.h defines int64_t as long under __CUDA_ARCH_RTC__, 32-bit on Win64. The <T,8> defect here is the same LLP64 class, which makes it a pattern in this codebase rather than a one-off. Both are worth fixing together, and both argue for a general "never spell a fixed-width type as long" rule.
• atomic.h: integer CAS specialisations retyped from upstream — signed char corrupts neighbouring bytes, unsigned long truncates on Win64 #98 (attribution) — the same file's provenance.
• Also stale and inert: upstream now guards atomicAddNoRet behind __builtin_amdgcn_processor_is("gfx908") and otherwise uses unsafeAtomicAdd; this copy retains the older unconditional #if defined(USE_ROCM) form. USE_ROCM is never defined in this project and atomic.h is the only file mentioning it, so this is vestigial rather than broken.

The #if __CUDA_ARCH__ < 600 double-precision CAS fallback and AtomicFPOp<double> were checked in the same pass and are byte-for-byte identical to current upstream, including the clang-CUDA enable_if hack and the integer-comparison loop that avoids a NaN hang. Those need no change.