#3172 gfx906 (2x MI50): build fails on 0.1.44+ — cublas* in qwen35, cuda.h in the nvfp4 MMQ target

本文

**Environment**

- Ubuntu 24.04.5, kernel `6.8.0-147-generic`, amdgpu from the kernel
- ROCm from the community image `mixa3607/rocm-gfx906:7.14-complete` (the one `docs/AMD_HIP.md` names for gfx906)
- 2× Instinct MI50 16 GB, reported by the driver as `1002:66af` (Vega 20), wave64
- Build inside that image, checkout at `477ac13` (main, 0.1.45):

```sh
cmake -S . -B build-906 -G Ninja \
  -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 \
  -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \
  -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++
ninja -C build-906 strata
```

**What happens**

Two independent failures, both from modules added after 0.1.43:

1. `src/qwen35/q36_engine.cu` — CUDA names the gfx906 compat layer does not map:

```
src/qwen35/q36_engine.cu:1555:13: error: use of undeclared identifier 'cublasSgemm'; did you mean 'hipblasSgemm'?
src/qwen35/q36_engine.cu:1958:29: error: use of undeclared identifier 'cublasGemmStridedBatchedEx'; did you mean 'hipblasGemmStridedBatchedEx'?
src/qwen35/q36_engine.cu:1972:29: error: use of undeclared identifier 'cublasGemmStridedBatchedEx'; did you mean 'hipblasGemmStridedBatchedEx'?
```

`CMakeLists.txt` adds `strata_qwen35` (and links it into `strata`) unconditionally in the CUDA branch, and `STRATA_HIP_GFX906` turns that branch into a HIP build, so the file is compiled with hipcc.

2. `src/prefill/mmq_nvfp4_w4a4.cu` — the target is built without `GGML_USE_HIP`, so llama.cpp's `common.cuh` falls through to the CUDA vendor header:

```
In file included from src/prefill/mmq_nvfp4_w4a4.cu:8:
In file included from .../ggml-cuda/common.cuh:40:
.../ggml-cuda/vendors/cuda.h:4:10: error: 'cuda.h' file not found with <angled> include
.../ggml-cuda/vendors/cuda.h:6:10: fatal error: 'cuda_bf16.h' file not found
```

Note that `strata_mmq` compiles its llama.cpp objects fine — only `strata_mmq_w4a4` misses the flag.

**Workaround that works**

`v0.1.43.1` (`a0116b2`) builds cleanly with the exact command above (that is what the machine now runs). So the regression is confined to 0.1.44+.

**Suggestions**

- Put the `qwen35` block behind `if(NOT STRATA_HIP_GFX906)`: `generate.cpp` already guards every reference with `#if defined(STRATA_HAVE_QWEN35)`, so a gfx906 build that does not define it links and runs. I verified this locally: with the block wrapped, the qwen35 error disappears and the build proceeds.
- Add `GGML_USE_HIP` (or the gfx906 compat include order) to the `strata_mmq_w4a4` target, or skip that target for `STRATA_HIP_GFX906`.
- If qwen35 is meant to work on gfx906 eventually, `cublasSgemm`/`cublasGemmStridedBatchedEx` need entries in `include/strata/platform/hip_compat/`.

Happy to test a patch on this machine — it is a dedicated 2×MI50 box with the community ROCm image installed.