[ROCm] Add AMD GPU support via ROCm/HIP build path

[jeffdaily]

This adds an AMD GPU backend to yalm built with ROCm/HIP, alongside the existing NVIDIA CUDA backend. The CUDA build is unchanged: with no extra flags the Makefile still drives nvcc exactly as before. Passing USE_HIP=1 instead compiles the .cu translation unit with hipcc and links the HIP runtime; the target architecture is selected with HIPARCH (default gfx90a, e.g. HIPARCH=gfx1100), so no source edit is needed to retarget.

Review order:

1. src/cuda_to_hip.h (new) — a compatibility shim. On the CUDA build it is a thin pass-through to the CUDA runtime headers, so nvcc behavior is untouched. On the HIP build it includes the HIP runtime and aliases only the CUDA runtime symbols yalm uses to their hip* equivalents. It is included by infer.cu (via hipcc) and by host C++ TUs (via model.h), so the host-side graph/state structs see the same type aliases.
2. Makefile — a USE_HIP=1 branch selecting hipcc, -lamdhip64, and --offload-arch=$(HIPARCH); host C++ flags gain -DUSE_HIP -D__HIP_PLATFORM_AMD__ and the link gets -no-pie. HIPARCH defaults to gfx90a and is caller-overridable; the CUDA branch is the original recipe, unchanged.
3. src/infer.cu — wave-width correctness, guarded so the CUDA path is unchanged:
   • matmul warps-per-block is derived at runtime from the device warp size (max-threads-per-block / warp_size): 16 on a 64-wide wavefront, 32 on a 32-wide one, keeping the block at the 1024-thread cap on every device. A compile-time arch macro cannot drive this, since the value is used in host launch code where device macros are undefined.
   • att_mix's __shared__ arrays are sized to the maximum wavefront width (64 on HIP, 32 on CUDA), since the kernel launches warpSize threads (64 on AMD).
   • The standalone kernel-test entry points query the device warp size and max-threads-per-block lazily rather than assuming 32 / 1024.
   • FULL_MASK is a 64-bit all-ones literal under HIP (ROCm's __shfl_*_sync require a 64-bit mask); CUDA keeps 0xffffffff.
   • Host-registered memory is handed to kernels via hipHostGetDevicePointer on HIP, with buffers registered using the mapped flag. Whether a host-registered pointer is directly device-addressable depends on the GPU's BAR aperture (large/resizable BAR), not the OS, so the device pointer is always resolved on HIP; on CUDA (UVA) this is a no-op.
4. src/model.h — include the shim instead of cuda_runtime_api.h directly.
5. README.md — ROCm/HIP build instructions next to the CUDA ones.

Test plan / validation

The bundled kernel test (make test USE_HIP=1 HIPARCH=<arch> && ./build/test) compares the GPU matmul / multi-head-attention / feedforward kernels against a CPU reference at epsilon 1e-4 and runs the CPU regression guard. Built and run on real AMD GPUs, deterministic across repeated runs:

• Linux gfx90a (Instinct MI250X, wave64), ROCm 7.2
• Linux gfx1100 (Radeon Pro W7800, RDNA3 wave32), ROCm 7.2
• Windows gfx1201 (Radeon RX 9070 XT, RDNA4) on ROCm

Both wavefront widths are covered, so the wave-width fixes are validated on real hardware.

On gfx1100, full end-to-end inference was additionally validated: Mistral-7B-Instruct-v0.2 (FP16) generates coherent text at ~26.7 tok/s, with the hipGraph capture/replay path confirmed exercised (hipGraphInstantiate / hipGraphLaunch succeed for prefill and decode), no NaN or GPU fault.

Authored with the assistance of Claude.