Pull requests / #3083

#3083 fix(HIP compat): the Qwen3.6 engine does not compile on ROCm, and gfx906 does not link at all

open · @benthic-io · 0 comentarios · En GitHub

BenchmarksAMD / HIPNVIDIA / CUDAModels & quants

Descripción

## Title

fix(HIP compat): the Qwen3.6 engine does not compile on a ROCm target

Issue: Resolves #

## Summary

Three separate defects, all on the AMD path, none reachable from a CUDA build:

1. `main` **does not link at all** on `STRATA_HIP_GFX906` - the Qwen3.6 / qwen35moe engine path cannot be built.
2. `src/qwen35/q36_engine.cu` does not compile for HIP - two cuBLAS names the HIP compatibility header does not map.
3. Two expert kernels in the shared file break any HIP build - a `__shared__ alignas(16)` the HIP host pass rejects.

None is a configuration problem. A `#define` cannot come from a CMake option or an environment variable, the
declaration syntax is fixed by the compiler, and a missing source file is missing.

## What changed

**`CMakeLists.txt`** - `STRATA_HIP_GFX906` forces `STRATA_ENABLE_CUDA ON`, so a gfx906 build takes this
`add_library(strata_mmq ...)` and not the `STRATA_PREFILL_MMQ` one below it - and this list omits
`src/prefill/moe_f16_kq.cu`, which the other one does list. `moe_fused_iq.cu` calls `f16kq::covers / gate_up /
down` inside `#if defined(__HIPCC__)` with no `STRATA_PF_F16` opt-in around them, so every HIP build on this path
references symbols that are not compiled:

```
ld.lld: error: undefined symbol: strata::prefill::fused::f16kq::covers(int)
ld.lld: error: undefined symbol: strata::prefill::fused::f16kq::gate_up(...)
ld.lld: error: undefined symbol: strata::prefill::fused::f16kq::covers_down(int)
ld.lld: error: undefined symbol: strata::prefill::fused::f16kq::down(...)
```

Adding the file resolves it; `f16kq::covers()` keeps the gfx11 kernels undispatched elsewhere, so run-time
behaviour is unchanged. Guarded by `STRATA_HIP_GFX906` because a CUDA build references none of these symbols and
should not compile a gfx11 kernel file.

**`include/strata/platform/hip_compat/cublas_v2.h`** - two missing cuBLAS-to-hipBLAS aliases:

- `cublasSgemm`, the F32 projections on the tiled split-K SGEMM (`q36_engine.cu:1555`)
- `cublasGemmStridedBatchedEx`, the prompt chunk's FP16 Q K^T and P V (`q36_engine.cu:1958`, `:1972`)

**`src/kernels/cuda/iq_kernels.cu`** - `STRATA_SHARED_INT4`, for the two `int4` staging arrays at `:1856` and
`:2043`. These spell the attribute out longhand where the neighbouring line uses `STRATA_SHARED_ALIGN16_U32`.
That macro exists for exactly this reason (#778) and its comment already records the workaround, so this follows
the established pattern. `int4` is 16-byte aligned by type, so under `__HIPCC__` the explicit attribute is
dropped - same layout, as the existing comment says for `uint32_t`.

## Extra Notes

Before, on gfx906:

```
q36_engine.cu:1555: error: use of undeclared identifier 'cublasSgemm'; did you mean 'hipblasSgemm'?
q36_engine.cu:1958: error: use of undeclared identifier 'cublasGemmStridedBatchedEx'; did you mean 'hipblasGemmStridedBatchedEx'?
iq_kernels.cu:1856: error: 'alignas' attribute cannot be applied to types
iq_kernels.cu:2043: error: 'alignas' attribute cannot be applied to types
... then, at link:
ld.lld: error: undefined symbol: strata::prefill::fused::f16kq::covers(int)
```

After: `ninja strata strata-q36` is clean.

Verified on **gfx906 (Radeon VII, 16 GiB), ROCm 7.1, Fedora 44**, gcc host compiler,
`-DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 -DSTRATA_GGML_DIR=<llama.cpp>`. Each file was checked
individually: it fails to compile without its change and compiles with it. The rebuilt `strata` then serves the
IQ4_XS pack (HTTP 200, correct answer, 3982 MiB free with everything loaded) and `strata-q36` runs a Qwen3.6
GGUF at 262144 cells (1329 tok/s pp2048, 49.3 tok/s tg32 with `--kv-q8`).

`src/prefill/gemm.cu` also names `cublasSgemm`, `cublasGemmAlgo_t` and `CUBLAS_STATUS_INTERNAL_ERROR`, but all
three are inside `#if defined(STRATA_EXPERIMENTAL_SM60)` (the experimental Pascal path) and are not compiled in
a normal HIP build, so no mapping is added for them here.

One thing I noticed but did **not** change, since it is a design question rather than a defect: the HIP
workarounds are gated on `STRATA_USE_HIP`, which `cmake/hip_backend.cmake:98` defines for the
`strata_hip_runtime` target - i.e. the `STRATA_ENABLE_HIP=ON` backend. `STRATA_HIP_GFX906=ON` instead defines
`STRATA_HIP_GFX906` and never `STRATA_USE_HIP`, so the 91 `#if defined(STRATA_USE_HIP)` sites across 12 files
are compiled out of a gfx906 build, including `STRATA_HIP_ADAPT_KERNEL_COPY=1` (#884, the SDMA hang workaround
in `generate.cpp:259`) and `STRATA_FILE_CACHE_OLD_RULE` (#1752, a gfx906-specific stall). Only
`generate.cpp`'s outer guard at `:248` is written as `(defined(STRATA_USE_HIP) || defined(STRATA_HIP_GFX906))`.
Worth deciding whether the two backends should agree here; happy to test a patch that defines `STRATA_USE_HIP`
on the gfx906 path if that is useful.

En el sitio

Enlaces a install, modelos, releases.