Pull requests / #3078

#3078 gfx906: compile the engine on a system ROCm 6.3.3 (blockers past the cublas maps)

open · @ken2425pm · 0 コメント · GitHub で見る

Server & APIMulti-GPUAMD / HIPNVIDIA / CUDADocumentation

本文

Complements #2460 / #2459 (the cublas name maps): with those in place, five more spots still block `-DSTRATA_HIP_GFX906=ON` on a **system** ROCm 6.3.3 (clang 18, `/opt/rocm-6.3.3`, no container). This PR carries only those, so the cublas PR keeps its scope.

**What it changes**

- `include/strata/platform/hip_compat/strata_hip.h`: void wrappers for the call sites that genuinely discard their `hipError_t` result — ROCm 6.x marks the type `[[nodiscard]]`, so every unchecked call warns under `-Wall -Wextra`. The pass-through wrappers are gone; those names map straight onto `hip*`.
- `src/kernels/cuda/iq_kernels.cu`: `__shared__ alignas(16) int4` does not compile on HIP (the #778 pattern) — a `STRATA_SHARED_ALIGN16_INT4` macro keeps the CUDA spelling and takes the attribute form under HIPCC.
- `src/qwen35/q36_engine.cu`: `__ldg` does not resolve in this TU under HIP (plain deref, semantically identical on gfx906); `__half_as_ushort` is device-only on HIP (memcpy bit-cast in the host debug print); `CK()` moves into a template so the void-returning compat wrappers type-check on both toolkits.
- `src/core/remote_experts.cpp`: `cudaInitDevice` has no HIP counterpart — excluded under `STRATA_HIP_GFX906` the way Metal already is.
- `src/kernels/qsa_prompt_attn_parity.cpp`: the sm_80+ Q4_0 gate must not query device attributes on a HIP build (the existing `__HIPCC__` exclusion, widened).

**What it leaves alone**: CUDA and SYCL are untouched — every hunk sits inside a HIP guard or a HIP-only macro branch, so those paths compile as `main` does by construction. No defaults change; no opt-in surface added.

**Verified** on 2x Instinct MI50 (gfx906), system ROCm 6.3.3: the `strata` target links clean with `-Wall -Wextra`, and generation runs (measurements in `docs/ROCm_6_3_3.md`, #3077).

**Not tested**: a CUDA build and SYCL (no CUDA/oneAPI toolchain on this box — stated, not skipped silently), and a `-Werror` build (`STRATA_WERROR=OFF`; the make-it-true path is measured in the report).

関連リンク

インストール・モデル・リリースへの站内リンク。