Pull requests / #1083
#1083 hip/gfx906: the tree does not build (three STRATA_USE_HIP gates miss STRATA_HIP_GFX906)
closed · @xxDoman · 0 commentaires · Sur GitHub
Multi-GPUAMD / HIPNVIDIA / CUDAModels & quants
Description
Re-creation of #1066 on the new main (the history rewrite closed it; the code is unchanged and the files are byte-identical to v0.1.40, so I cherry-picked my commits onto `main`). ## Problem Building for gfx906 (`-DSTRATA_HIP_GFX906=ON`) fails with three compile errors. The gfx906 build defines `STRATA_HIP_GFX906`, **not** `STRATA_USE_HIP` (that one belongs to the wave32 HIP backend). Three sites gate only on `STRATA_USE_HIP`, so on gfx906 they keep the CUDA branch and do not compile. #808 fixed the same class of thing in `strata_hip.h`. ## Errors and fixes **1. `src/kernels/cuda/fused_gr.cu`** — a bare `return;` in the `bool fused_gr_read_multi`, inside the `STRATA_HIP_GFX906` `STRATA_GR_SPLIT` block: ``` fused_gr.cu:1743:9: error: non-void function 'fused_gr_read_multi' should return a value [-Wreturn-mismatch] ``` → `return false;` — the q8_1 contract of the other multi-token paths (no q8_1; the caller quantizes). The v3 path above it already returns `false`. **2. `src/core/mtp.cpp`** — the MTP stream-priority block: ``` mtp.cpp:421:33: error: 'cudaDeviceGetStreamPriorityRange' was not declared in this scope mtp.cpp:422:16: error: 'cudaStreamCreateWithPriority' was not declared in this scope ``` → gate the block on `!defined(STRATA_HIP_GFX906)` as well (HIP has no stream priority; the default priority is the intended behaviour, per the file's own comment). **3. `src/core/vmm.cpp`** (new in 0.1.40) — includes `<cuda.h>` for the CUDA driver VMM API: ``` vmm.cpp:4:10: fatal error: cuda.h: No such file or directory ``` → the same gate. The file already has the HIP fallback in its `#else` branch. ## Scope / safety Every affected line is inside an `#if` only the gfx906 build reaches, so CUDA and the wave32 RDNA backend are untouched. ## Verified ``` cmake -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 -DCMAKE_HIP_COMPILER=/opt/rocm/lib/llvm/bin/clang++ -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_BUILD_TYPE=Release .. ``` `[100%] Linking HIP executable strata`, EXIT 0, links only `libhipblas/libamdhip64/librocblas/libhipblaslt` (ROCm 7.2.1). The engine starts on a 2x gfx906 rig (Instinct MI50 32 GB + Radeon VII 16 GB) and generates. The `dp4a` item is a separate PR (parity, not build-blocking).
Sur le site
Liens install, modèles, releases.