Pull requests / #1066

#1066 hip/gfx906: the 0.1.40 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

## Problem

Building 0.1.40 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 already 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 three lines up 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, as the file's own comment says).

**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 (`vmm_available() == false`).

## Scope / safety

Every affected line is inside an `#if` that only the gfx906 build reaches, so CUDA and the wave32 RDNA backend are untouched. No behaviour change on any other backend.

## 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.

Separate from this: a performance item (gfx906 has `v_dot4_i32_i8` but `dp4a` falls back to a 4-iteration loop) — I will open that on its own, since it needs a parity test rather than being build-blocking.

Sur le site

Liens install, modèles, releases.