改动 / #3145

#3145 fix(HIP): pair gfx1031 with gfx1030 and run it as 10.3.0 - ROCm 7.1.1's rocBLAS ships no gfx1031 kernels

open · @ljmbot · 0 评论 · 去 GitHub 看

BenchmarksSetup & installAMD / HIPNVIDIA / CUDAModels & quantsDocumentationWindows

说明

Issue: Resolves #

## Summary
On an RX 6700 XT (`gfx1031`) the engine did not compile, and once it compiles it loads, answers a 46-token prompt at
20-27 tok/s, and dies on the first prompt of about 12K tokens. The death message looks like an out-of-memory:

```
rocBLAS error: Cannot read /opt/rocm/lib/rocblas/library/TensileLibrary.dat: No such file or directory for GPU arch : gfx1031
```

It is not memory: 62.7 GiB RAM, 12,272 MiB VRAM, 4.5 GiB VRAM free after the load. The prompt path is the only path
that calls rocBLAS (decode streams experts from host RAM and uses the CPU gemv), and ROCm 7.1.1's rocBLAS ships 31
`gfx1030` library files and 0 `gfx1031` ones (`ls /opt/rocm/lib/rocblas/library | grep -c gfx1031` is 0). So the
engine built `gfx1030;gfx1031` fixes the kernels and leaves the library unfixable, and
`HSA_OVERRIDE_GFX_VERSION=10.3.0` is what makes rocBLAS find that library at runtime. This PR makes setup.py do both,
adds a no-GPU test for it, fixes the compile, and documents it. No new default; a card whose arch has a Tensile
library gets exactly what it got before.

Measured on: RX 6700 XT (12,272 MiB), Xeon E5-1620 v3, 64 GB RAM, Ubuntu 24.04.5, kernel 6.8.0-147, ROCm 7.1.1 in
/opt/rocm, rocBLAS 5.1.1, Strata 0.1.45 at commit `477ac13c`, Swift 1.5 IQ3_S, 65,536-token context, KV int8,
`--expert-cache auto`, `--spec 4`, engine built from source on this machine.

## What changed
- `setup.py` (43 added, 4 changed): `HIP_ROCBLAS_SIBLINGS = {"gfx1031": "gfx1030"}` and
  `HIP_ROCBLAS_AS = {"gfx1031": "10.3.0"}` next to `AMD_ARCHS`, with the reason and these numbers in the comment
  above them. `HIP_ARCH` falls back to the sibling when the card's own arch has no Tensile library, so a source build
  gets the gfx1030 kernels, and the run config's `env` gets `HSA_OVERRIDE_GFX_VERSION=10.3.0`, which is what makes
  rocBLAS find that library at runtime. The env entry only appears for a card whose arch has no library. `HIP_ARCH`
  and the arch list are the two points where a fat build becomes necessary; the prebuilt `strata-win-hip.zip` has the
  same hole, so the release binary alone does not cover the card.
- `tools/test_setup_gfx1031.py` (new, 77 lines): a 6700 XT gets `["gfx1030", "gfx1031"]`, the override in the config
  and in the printed export line; a 9070 XT is untouched; a 9070 XT keeps `12.0.0`. Runs without a GPU and without
  downloads.
- `src/kernels/cuda/iq_kernels.cu` (4 added, 2 changed): `__shared__ alignas(16) int4 s_u[...]` in
  `expdec_gu_kernel` (1858) and `expdec_down_kernel` (2045) does not compile under HIP clang - "'alignas' attribute
  cannot be applied to types", the #778 pattern - so the engine did not build on this machine. The file already
  carries the answer for `uint32_t` shared arrays: `STRATA_SHARED_ALIGN16_U32` puts `__attribute__((aligned(16)))`
  after the declarator under `__HIPCC__` and keeps `alignas` elsewhere. This adds the same macro for `int4`
  (`STRATA_SHARED_ALIGN16_I4`) and uses it at those two sites. Layout identical, and nvcc accepts the CUDA form,
  which is why CUDA builds never see this.
- `docs/AMD_GFX1031.md` (new, 88 lines): the symptom, the two changes, the probe numbers, what to copy for this card,
  and the `STRATA_HIP_PROMPT_F16` measurement. `docs/AMD_HIP.md` (9 added, 1 changed): a line pointing at it for
  cards whose arch ships no rocBLAS kernels.

How I checked:
- `python3 tools/test_setup_gfx1031.py` -> `Ran 8 tests ... OK`; `tools/test_setup_choices.py` -> `Ran 47 tests ...
  OK`; `tools/test_setup_amd.py` -> `Ran 40 tests ... OK`. No GPU, no downloads.
- the engine built from source here: `engine/BUILD.json` `"archs": ["gfx1030", "gfx1031"]`, `engine/strata`
  53,696,584 bytes (53.7 MB) against 32.5 MB for one arch, 275 s to build on this 8-thread CPU, and no run-time cost.
- the built-in suite in two arms (`tools/strata_bench.py --long`, three runs per test, in the paired report PR
  #3146): prompt 339.2 (294.1-339.5) / 352.5 / 356.2 tok/s at 4K / 16K / 32K becoming 595.9 / 692.9 / 707.5 with
  `STRATA_HIP_PROMPT_F16=1` (1.76x / 1.97x / 1.99x), TTFT 11.90 / 45.61 / 90.05 s becoming 6.80 / 23.23 / 45.39 s,
  chat decode 21.7 -> 21.0 tok/s over 9 runs each (decode does not use these GEMMs).
- the same 11,103-token prompt in one server session: 33.9 s / 32.8 s / 34.4 s reading (327.6 / 338.1 / 322.4 tok/s)
  on the default route and 17.9 s (620.9 tok/s, 1.90x) with the switch; a 13,023-token prompt 38.6 s (337.6 tok/s);
  the second copy of the 11,103-token prompt reuses 11,098 of 11,103 tokens in 102 ms.
- `scripts/rocblas_gemm_probe.cpp` (in the report PR) at `N = 10240 x T = 1024 x K = 2560`, each type checked against
  a CPU reference: 7.0 TFLOPS f16 in / f32 out (the default route), 21.3 f16 in / f16 out, 6.6 bf16, 11.8 f32.
- the two halves are both needed: the override on a `gfx1031`-only binary SEGFAULTS (HIP then looks for gfx1030 code
  objects) and the fat binary alone leaves rocBLAS asking for gfx1031. A `ROCBLAS_TENSILE_LIBPATH` directory of
  symlinks renaming the gfx1030 files to gfx1031 is not a fix - the abort stops and every `rocblas_gemm_ex` returns
  status 6 (internal_error) and writes nothing.

## Extra Notes
What this leaves alone: no new default and no behaviour change for a card whose arch has a Tensile library (gfx1151,
gfx1201, gfx942, RDNA2, RDNA1 - gfx1151 already pairs with gfx1151, so it takes the same code path and gets the same
arch back); the CUDA path is untouched (the non-`__HIPCC__` branch of the new macro is the CUDA form, as it was); the
`iq_kernels.cu` change is a compile fix with identical layout, not a new code path; no other engine source is touched.

What is not measured here: there is no built-in-suite arm for a `gfx1031`-only build - that build dies on the first
long prompt, so the only before/after for the arch change and the override is the 11,103-token pair above and the
probe. Also not tested: 4K / 32K token-id parity between the two GEMM routes, the other `gfx1031` cards (6750 XT,
6800M), Windows, TheRock's venv ROCm, a layer split with a 6700 XT in it, ROCm older than 7.1.1, PCIe width, RAM
speed, temperatures.

本站相关内容

相关页面的快捷入口。