[ROCm] Add portable HIP build for ZhiLight on AMD GPUs#81
Open
jeffdaily wants to merge 3 commits into
Open
Conversation
Adds a ROCm/HIP build path to ZhiLight (CUDA-only before this). The portable inference surface -- elementwise/norm/softmax/rope/embedding, the scalar attention decode path, the cublasLt GEMM, the int8 reduce kernels, and the fp16/bf16 MoE expert-routing softmax -- now builds with hipcc and runs correctly on gfx90a (CDNA2, wave64) and gfx1100 (RDNA3, wave32). The NVIDIA tensor-core / Hopper-only fast paths are gated OFF on this pass and left as a deferred follow-up (see "Deferred" below). The NVIDIA dispatch and the deferred fast paths are unchanged; the only shared (non-USE_HIP-guarded) edit is the block-reduce partial-selection guard, rewritten in a BM_WARP_SIZE-driven form that is value-identical on CUDA (BM_WARP_SIZE = 32, ceil == floor for the 32-multiple block sizes used here). The int8 quant reduce and the MoE routing softmax use the width-32 logical-warp reduction (warpReduceMaxWidthB / warpReduceSumWidthB) already used by attention_kernel.cu, all under USE_HIP -- equivalent to the old full-warp reduce for a 32-element group on a 32-thread CUDA warp, and correct (no cross-row / inactive-lane read) on wave64. Everything else is behind USE_HIP / __HIP_PLATFORM_AMD__. This work was authored with assistance from Claude (Anthropic). Approach: a compat header plus CMake HIP language gating. README.md documents the ROCm build alongside the existing CUDA instructions. Suggested review order: 1. CMakeLists.txt + 3rd/bmengine/.../CMakeLists.txt: an opt-in USE_HIP that enable_language(HIP), reads CMAKE_HIP_ARCHITECTURES (default gfx90a only when unset, so other GPUs build gfx1100/gfx1151 with no source edit), marks the existing .cu LANGUAGE HIP, swaps cuBLAS/cuBLASLt/cuRAND/NCCL -> hipBLAS/hipBLASLt/hipRAND/RCCL, makes the flash-attn .so non-fatal, gates the deep_gemm/flash_mla H20 static libs off, pins -ffp-contract=on, and disables pybind11 IPO/LTO (the HIP link leaves a PyInit-less slim .so otherwise). 2. hip_compat/: the single compat header plus toolkit-named forwarding shims (cuda_runtime.h, cublasLt.h, nccl.h, ...) placed on the HIP-only include path so host .cpp resolve the CUDA spellings to HIP without editing include lines. 3. The wave64 root cause -- 3rd/bmengine/.../functions/reduce.cuh: warpReduce* used 32-lane offsets with 0xFFFFFFFF masks and blockReduce* hardcoded %32, /32 and shared[33]. On a 64-lane wavefront this silently reduced only half the lanes (wrong amax/softmax/etc., no crash). Rewritten wave-agnostically (BM_WARP_SIZE = 64 on __GFX9__ else 32; full-width no-mask shuffles on HIP; num_warps-driven block reduce). The block-reduce partial-selection guard is USE_HIP-gated so the CUDA path keeps its original spelling. This header is reused everywhere, so it is the highest-leverage edit. round_up_thread() now rounds the block size up to a whole wavefront (64 on HIP) so a reduction never runs on a half-filled wavefront. attention_kernel.cu keeps its 128-dim-head=lane*4 tiling as width-32 logical-warp reductions (correct on wave32 and wave64); the bitonic sorts in topk.cu/ff_kernel.cu use the no-mask wavefront shuffle. The new warpReduceSumWidthB<T,W> here is the broadcast-sum counterpart of the existing warpReduceMaxWidthB (width-W __shfl_down reduce + width-W __shfl(.,0) broadcast). 4. quant_reduce_kernel.cu and ff_kernel.cu DEV_softmax_inplace: two kernels reduce one fixed 32-element group but launch a 32-thread block or dim3(32,32), i.e. a half-filled (or two-row) wave64 wavefront. The old full-wavefront warpReduceMaxB / warpReduceSumB then fold the neighbouring row or inactive lanes 32-63 into the result on wave64. Both now use the width-32 form. ff_kernel.cu DEV_softmax_inplace is the fp16/bf16 MoE expert-routing softmax (KERNEL_top_k_softmax launches <<<gridDim, 32>>>): feedforward.cpp routes every fp16/bf16 MoE model through the generic MOEImpl -> route() -> top_k_softmax, so this is a SUPPORTED gfx90a path, not a deferred one (the standard Mixtral / Qwen-MoE softmax routing when topk_group <= 1). The KERNEL_group_topk softmax (topk_group > 1) launches num_group*WARP_SIZE >= 64 threads and stays on the wave-correct block-reduce branch, so it is untouched. 5. The remaining files are the library swaps and the gating of the deferred NVIDIA-only paths. Deferred (NVIDIA tensor-core / Hopper-only, reimplement-not-port; throw a clear error or fall back on HIP): marlin / awq / fp8 / gptq quant kernels (mma.sync, ldmatrix, cp.async, cvt.e4m3x2, prmt.b32, lop3, __vsub4), the wmma attention decode path, int8-quantized-KV attention, the deep_gemm FP8-block + flash_mla MLA H20 static libs (DeepSeek-V3 FP8-block + MLA fast paths), and flash-attn prefill. The unquantized fp16/bf16 path (dense + MoE), the fp16/bf16 MoE routing softmax, and the cublasLt Int8Linear GEMM are supported (the int8 tensor-parallel all-reduce quant kernels are now wave64-correct but unexercised by the kernel-test gate); MLA falls back to the generic attention path. GPTQ/AWQ/Marlin/fp8 quantized MoE remain deferred. Test Plan Built for gfx90a and ran the in-tree GPU kernel tests (torch reference, no model weights) on a real MI250X GCD; all pass and are deterministic across runs: ``` HIPCXX=/opt/rocm/lib/llvm/bin/clang++ ZHILIGHT_USE_HIP=1 ENABLE_NCCL_TP=on TESTING=1 \ cmake -S . -B build -GNinja -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_HIP_PLATFORM=amd \ -DCMAKE_BUILD_TYPE=Release -DWITH_TESTING=ON -DPYTHON_EXECUTABLE=$(which python3) cmake --build build --target C internals_ -j16 HIP_VISIBLE_DEVICES=1 PYTHONPATH=. python3 -m pytest \ tests/test_softmax.py tests/test_attn_softmax.py tests/test_rotary_embedding.py \ tests/test_feedforward.py tests/test_linear.py tests/test_embedding.py \ tests/test_arthmetic.py tests/test_concat_tensor.py tests/test_index_along_dim.py \ tests/test_log_prob.py tests/test_attention.py tests/test_lazy_loader.py # 77 passed, deterministic across runs ``` The kernel-test gate does not cover the MoE routing softmax, so the DEV_softmax_inplace fix was validated separately by binding nn::top_k_softmax into the internals test module and comparing the routing softmax + top-k against a PyTorch reference over standard Mixtral / Qwen-MoE shapes (fp16 and bf16, 8 to 128 experts, top-2/6/8): max value error 1.2e-7, expert-id sets exact, and bit-identical across re-runs (no inactive-lane reads). AMD_LOG_LEVEL=3 confirms KERNEL_top_k_softmax<__half> dispatches on amdgcn gfx90a. That binding is a throwaway validation aid and is not committed. Toolchain: ROCm 7.2.1, ROCm PyTorch 2.13. Also built and validated on gfx1100 (RDNA3, wave32). The amax/amin tests fail before the wave64 reduce + round_up_thread fix and pass after, confirming the root cause.
The pin set CMAKE_HIP_ARCHITECTURES to gfx90a before enable_language(HIP) whenever it was unset, which preempted enable_language(HIP)'s own host-GPU detection. A user on a non-gfx90a AMD GPU who omitted -DCMAKE_HIP_ARCHITECTURES silently built gfx90a code objects that fail to load on their card at runtime. Removing the pin relies on enable_language(HIP): it honors an explicit -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the host GPU(s), and errors out when no GPU is found (so a no-GPU build host must set the arch explicitly rather than getting a silently wrong default). Explicit-arch builds are unchanged. Authored with assistance from Claude.
The bundled bmengine CMake pinned CMAKE_HIP_ARCHITECTURES to gfx90a before
enable_language(HIP) when it was unset, which preempted CMake's own host-GPU
detection: a user on a non-gfx90a AMD GPU who did not pass
-DCMAKE_HIP_ARCHITECTURES would silently build gfx90a code objects that fail to
load on their card ("no kernel image"). Removing the pin lets enable_language(HIP)
honor an explicit -DCMAKE_HIP_ARCHITECTURES, auto-detect the host GPU(s), or error
on a no-GPU build host. Mirrors the same fix already applied to the top-level
ZhiLight CMake.
Test Plan: build is byte-identical for the explicit-arch builds used in
validation (cmake -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a ...); the change
only affects the unset-default path.
Authored with the Claude AI assistant.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds a ROCm/HIP build path to ZhiLight, which was CUDA-only before this. The portable inference surface -- elementwise/norm/softmax/rope/embedding, the scalar attention decode path, the cuBLASLt GEMM, the int8 reduce kernels, and the fp16/bf16 MoE expert-routing softmax -- now builds with hipcc and runs correctly on AMD GPUs. Everything is gated behind an opt-in
USE_HIP; the NVIDIA build path is unchanged.Build it with:
Approach
A central compat header (
hip_compat/) plus toolkit-named forwarding shims (cuda_runtime.h,cublasLt.h,nccl.h, ...) on a HIP-only include path, so host.cppresolve the CUDA spellings to their HIP equivalents without editing include lines. CMake gains an opt-inUSE_HIPthat enables the HIP language, marks the existing.cuas HIP, swaps cuBLAS/cuBLASLt/cuRAND/NCCL for hipBLAS/hipBLASLt/hipRAND/RCCL, and defaultsCMAKE_HIP_ARCHITECTUREStogfx90a(overridable for other GPUs with no source edit).The substantive correctness work is wave64:
3rd/bmengine's warp/block reductions hardcoded a 32-lane warp (32-lane shuffle offsets,0xFFFFFFFFmasks,%32//32,shared[33]). On a 64-lane wavefront these silently reduced only half the lanes (wrong amax/softmax, no crash). They are rewritten wave-agnostically underUSE_HIP; the CUDA path keeps its original spelling and is value-identical. The README documents the ROCm build alongside the existing CUDA instructions. The commit body has a file-by-file review order.Scope
Validated on real GPUs on Linux: gfx90a (MI250X, CDNA2/wave64) and gfx1100 (RDNA3/wave32). Windows is not part of this PR -- ZhiLight's bundled
bmenginehost runtime uses Linux-only POSIX APIs (execinfo.h,sched.h,pthread.h,sys/syscall.h) that have no Windows equivalent; the GPU kernels themselves compile for AMD on Windows, but the host runtime would need a separate Windows-compat effort.The NVIDIA tensor-core / Hopper-only fast paths (Marlin/AWQ/FP8/GPTQ quant kernels, the WMMA attention decode, int8-quantized-KV attention, the deep_gemm FP8-block and flash_mla MLA H20 static libs, and flash-attn prefill) are gated off on HIP and throw a clear error or fall back. The unquantized fp16/bf16 dense + MoE path, the fp16/bf16 MoE routing softmax, and the cuBLASLt Int8Linear GEMM are supported; MLA falls back to the generic attention path.
Test plan
Built for gfx90a and ran the in-tree GPU kernel tests (torch reference, no model weights) on a real MI250X GCD; all pass and are deterministic across runs:
The amax/amin tests fail before the wave64 reduce fix and pass after, confirming the root cause. Also built and validated on gfx1100. Toolchain: ROCm 7.2.1, ROCm PyTorch 2.13.
This work was authored with assistance from Claude (Anthropic).