Issues / #2459
#2459 gfx906/HIP build broken on v0.1.44.1: two compile blockers (platform/hip_compat/cublas_v2.h missing three cublas maps, and the new strata_mmq_w4a4 target missing GGML_USE_HIP)
open · @kietq8684-design · 0 Kommentare · Auf GitHub
AMD / HIPNVIDIA / CUDADocumentation
Beschreibung
### Summary
`-DSTRATA_HIP_GFX906=ON` no longer builds on the latest tag. v0.1.42 (`61b3fb5`) builds clean; v0.1.44.1 (`296b59d`)
fails with two independent blockers, both introduced after 0.1.42. `docs/AMD_HIP.md:472` says the gfx906 engine
"is checked to compile for gfx906", so this is a regression rather than a missing feature. A 9-line patch makes
`ninja -C build-906 strata` finish with 149/149 linked, exit 0.
### Environment
- GPU: Radeon VII 16 GB (`0x66af`) + Radeon Pro VII / Instinct MI50 16 GB (`0x66a1`), both gfx906
- ROCm: HIP 7.14 (hipcc 7.14.60850) from `mixa3607/rocm-gfx906:7.14-complete`
- Kernel: 6.8.0-146-p2p, Ubuntu 24.04, `HSA_OVERRIDE_GFX_VERSION=9.0.6`
- Build: `-DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 -DSTRATA_GGML_DIR=/llama
-DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++`, Ninja
- llama.cpp at the pinned commit (`3cf03257f`)
- `git describe` on the tree: v0.1.42 (`61b3fb5`) builds; worktree at v0.1.44.1 (`296b59d`) fails
### Blocker 1 — `include/strata/platform/hip_compat/cublas_v2.h` is missing three mappings
`include/strata/hip_compat/cublas_v2.h` gained `cublasSgemm`, `cublasSgemmStridedBatched` and
`cublasGemmStridedBatchedEx` between 0.1.42 and 0.1.44.1, and `src/qwen35/q36_engine.cu` (new in 0.1.44.1) uses
them. But that file is **not on the gfx906 include path**: `CMakeLists.txt:60` says the cuBLAS names are mapped by
`include/strata/platform/hip_compat/`, lines 138/141/142 attach it `BEFORE INTERFACE`, and the only strata-side
`-I` in `compile_commands.json` is `-I/src/include/strata/platform/hip_compat`. That second file still maps
`cublasGemmEx` only, so the three names are undeclared.
Fix — add to `include/strata/platform/hip_compat/cublas_v2.h`:
```diff
#define cublasGemmEx hipblasGemmEx
+#define cublasSgemm hipblasSgemm
+#define cublasSgemmStridedBatched hipblasSgemmStridedBatched
+#define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx
```
### Blocker 2 — the new `strata_mmq_w4a4` target has no `GGML_USE_HIP`
```
FAILED: CMakeFiles/strata_mmq_w4a4.dir/src/prefill/mmq_nvfp4_w4a4.cu.o
/llama/ggml/src/ggml-cuda/vendors/cuda.h:4:10: error: 'cuda.h' file not found with <angled> include
/llama/ggml/src/ggml-cuda/vendors/cuda.h:6:10: fatal error: 'cuda_bf16.h' file not found
```
`mmq_nvfp4_w4a4.cu` includes llama.cpp's `common.cuh`, which selects `vendors/hip.h` only under
`#if defined(GGML_USE_HIP)` and otherwise pulls `vendors/cuda.h` (`<cuda.h>`, `<cuda_bf16.h>` — absent in a HIP
environment; `platform/hip_compat/` provides only `cuda_fp16.h`, `cuda_runtime.h`, `math_constants.h`).
`CMakeLists.txt:1791` sets `GGML_USE_HIP` on `strata_mmq` for the gfx906 branch and `:1576` has the `=1` form for
gfx11, but the `strata_mmq_w4a4` target added at `:1449-1457` was missed. Its sibling `mmq_nvfp4_w4a8.cu` has the
define and compiles (877 KB `.o`), which is what identifies this as an omission rather than a design choice.
CUDA builds are unaffected — they want `vendors/cuda.h`.
Fix — in `CMakeLists.txt`, after `target_link_libraries(strata_mmq PRIVATE strata_mmq_w4a4)`:
```diff
target_link_libraries(strata_mmq PRIVATE strata_mmq_w4a4)
+ # gfx906 / HIP: this unit includes ggml-cuda's common.cuh, which only takes vendors/hip.h when
+ # GGML_USE_HIP is defined (strata_mmq gets it in the HIP branches below). Without it the CUDA
+ # vendor headers (<cuda.h>, <cuda_bf16.h>) are pulled in and the HIP build fails.
+ if(STRATA_HIP_GFX906 OR STRATA_ENABLE_HIP)
+ target_compile_definitions(strata_mmq_w4a4 PRIVATE GGML_USE_HIP)
+ endif()
```
### Result with the patch
`ninja -C build-906 strata` → **149/149 linked, exit 0**, `build-906/strata` = 26,224,152 B (v0.1.42 is 21 MB).
Patch total: 9 lines (`cublas_v2.h` +3, `CMakeLists.txt` +6). The 0.1.42 tree is untouched.
### Note (not a blocker, just a build-time observation)
`strata_mmq_w4a4` is a unit that the gfx906/HIP path does not use at all, so an alternative to the define above
is to skip the target on this branch entirely — either way the HIP build should not be pulling CUDA vendor headers
for it.
Mehr auf der Site
Links zu Install, Modellen, Releases.