[ROCm] Add a HIP/ROCm build for AMD GPUs

[jeffdaily]

Brings up llm.c on ROCm with a minimal, additive port: one compat header (llmc/cuda_to_hip.h) force-included on the HIP translation units, a set of forwarding shims under llmc/hip_shims/ for the CUDA-named toolkit headers, and a USE_HIP=1 branch in the Makefile that compiles the existing .cu sources with hipcc and keeps the upstream target names. The NVIDIA build is unchanged: every source edit is behind a USE_HIP guard or a macro/cast that resolves to the original spelling on CUDA.

Authored with the assistance of Claude, an AI assistant by Anthropic.

Review order: start with llmc/cuda_to_hip.h (the only file that knows about HIP: runtime/bf16/cuBLAS(Lt) aliases, the 64-bit full-warp mask, streaming load/store and cooperative-groups reduce shims HIP lacks) and the Makefile HIP branch, then the arch-unified fault-class fixes in llmc/{cuda_common,cuda_utils, matmul,layernorm}.* and train_gpt2_fp32.cu.

Wavefront width is configurable, never a literal: AMDGPU_TARGETS drives LLMC_WARP_SIZE (64 on CDNA gfx9xx, 32 on RDNA), defined for both hipcc compile passes so the host launch geometry and the device reductions agree. The target arch is auto-detected with amdgpu-arch (mirroring the existing nvidia-smi compute_cap query on the CUDA path), with gfx90a as the fallback and AMDGPU_TARGETS= as an override, so other AMD GPUs (gfx1100/gfx1151) build with no flags or source change.

Fault-class fixes, all arch-unified (correct on wave32 and wave64):

• 64-bit _shfl*_sync masks (the 0xFFFFFFFF literal does not compile on ROCm);
• warpReduceSum/Max start offset = WARP_SIZE/2 so a wave64 reduction folds all 64 lanes, not 32;
• blockReduce cross-warp scratch sized to a compile-time upper bound;
• matmul_backward_bias grid stride and launch dims keyed on the wavefront width (a literal 32 read past the buffer / tripped a kernel assert on wave64);
• layernorm_backward_kernel10 temp-shared reservation matched to the kernel's WARP_SIZE-based offsets (a literal 32 over-allocates the dynamic shared mem on wave64; the substitution makes the host reservation exact, not a corruption fix -- the device tmp offsets already use the device WARP_SIZE);
• hipBLASLt has no MATMUL_DESC_SCALE_TYPE attribute (guarded out; the scale type follows the FP32 compute type);
• streaming __ldcs/__stcs/__stcg and cg::reduce/plus/greater shimmed (HIP lacks __stcs/__stcg and provides __ldcs only for half/half2; CG lacks all three);
• clang's -ffast-math is dropped in favor of -ffp-contract=fast -fno-math-errno: -ffast-math reassociates the softmax/layernorm-backward reductions into NaN gradients on gfx90a (the forward loss stays correct), which nvcc's --use_fast_math does not.

cuDNN flash attention stays off (the in-repo attention.cuh is the default); NCCL->RCCL is wired but exercised single-GPU. A determinism shift vs CUDA in BF16 is expected (the WARP_SIZE-seeded stochastic-rounding RNG and the wave64 accumulation order differ), so BF16 is validated by tolerance and loss-curve, not bitwise.

Test Plan (AMD Instinct MI250X, gfx90a, ROCm 7.2.1, one isolated GCD):

cd llm.c
./dev/download_starter_pack.sh   # PyTorch reference .bin files

make USE_HIP=1 AMDGPU_TARGETS=gfx90a NO_MULTI_GPU=1 NO_USE_MPI=1 -j 16 \
     test_gpt2fp32cu train_gpt2fp32cu test_gpt2cu train_gpt2cu

HIP_VISIBLE_DEVICES=2 ./test_gpt2fp32cu          # overall okay: 1
HIP_VISIBLE_DEVICES=2 ./test_gpt2cu              # forward + 14/16 grads OK;
HIP_VISIBLE_DEVICES=2 ./test_gpt2cu -r 0         # ln1b/lnfb exceed the
HIP_VISIBLE_DEVICES=2 ./test_gpt2cu -r 2         # NVIDIA-tuned BF16 thresholds
HIP_VISIBLE_DEVICES=2 ./test_gpt2cu -w 0         # on one near-zero element
HIP_VISIBLE_DEVICES=2 ./test_gpt2cu -b 32        # (expected determinism shift)

PRECISION=FP32 make USE_HIP=1 AMDGPU_TARGETS=gfx90a NO_MULTI_GPU=1 NO_USE_MPI=1 train_gpt2cu
HIP_VISIBLE_DEVICES=2 ./train_gpt2cu -b 1 -t 64 -d 256 -l 0.0001 -v 200 -s 200 \
     -a 1 -x 10 -r 0 -f 0 -e gpt2_124M.bin > out.txt
python dev/loss_checker_ci.py -f out.txt -s 20 -e 28 -a 5.0   # Success (<0.05%)

HIP_VISIBLE_DEVICES=2 OMP_NUM_THREADS=8 ./train_gpt2cu        # tinyshakespeare:
   # train loss 4.29 -> 3.38, val loss 4.51 -> 3.49, coherent text, no NaN

# CPU regression (port leaves the C path unchanged)
make train_gpt2 test_gpt2 && OMP_NUM_THREADS=8 ./test_gpt2   # overall okay: 1
cd dev/test && make PRECISION=BF16 test_dataloader && ./test_dataloader  # all OK

Windows (gfx1151) build enablement + a newer-ROCm fix:

• Makefile: the USE_HIP toolchain branch lived only under ifneq($(OS),Windows_NT), so Windows never selected hipcc. Added the mirror HIP branch to the Windows section (cmd where for hipcc detection + $(firstword); HIP flags; -DNOMINMAX/-DWIN32_LEAN_AND_MEAN because the HIP runtime headers pull in <windows.h> whose min/max macros break std::min; hipblaslt ships only libhipblaslt.dll.a on Windows (no .lib), linked by full path via -Xlinker; HIP_LIB_DIR passes the ROCm lib dir; .exe output handling for clang). Build with mingw32-make SHELL=cmd.exe USE_HIP=1 AMDGPU_TARGETS=gfx1151 HIP_LIB_DIR=<rocm>/lib ....
• cuda_to_hip.h: the cg::plus/greater/reduce fallback (for ROCm 7.2.x, which lacks cg::reduce) is now guarded #if HIP_VERSION < 71300000 -- newer ROCm (>= 7.13) ships them natively and redefining is an error. No-op on the 7.2.x lead.

Builds clean on Windows gfx1151 (test_gpt2cu.exe + train_gpt2cu.exe).

Validated on gfx1151 (Radeon 8060S APU, Windows, TheRock ROCm): test_gpt2fp32cu -> "overall okay: 1" -- GPT-2 FP32 training fwd+bwd loss matches the PyTorch reference at every step (5.27->0.38 over 10 steps, ref-matched to ~1e-4); exercises attention, layernorm, matmul (rocBLAS), cooperative-groups reductions and AdamW on wave32. The cg fix above splits the version guard: ROCm 7.13 ships cg::plus/cg::greater (so they are guarded out) but still has NO cg::reduce, so reduce is always supplied. The BF16 driver (test_gpt2cu) builds but its hipBLASLt GEMM crashes inside libhipblaslt.dll on gfx1151 (0xC0000005 once the Tensile kernels load) -- a TheRock hipBLASLt/gfx1151 runtime issue, NOT a port defect (the same kernels validate via the FP32/rocBLAS path).