JeanP00l 9ff290b493 hip: AMD gfx906 (Instinct MI50 / MI60, Radeon VII; wave64) as an opt-in build, STRATA_HIP_GFX906=ON ...

The HIP backend (STRATA_ENABLE_HIP) is wave32 RDNA only, and the engine refuses a wave64 card. This adds a
separate opt-in build for gfx906: the CUDA sources compiled as HIP through a small compat layer, with a CUDA warp
as a logical half of the 64-lane wavefront. With the option off nothing changes for CUDA or the RDNA backend;
every new code path is under STRATA_HIP_GFX906.

Build: -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 (implies STRATA_ENABLE_CUDA's targets, refuses
STRATA_ENABLE_HIP). The CUDA runtime/cuBLAS names come from include/strata/platform/hip_compat/ (first on the
include path), CUDA::cudart / CUDA::cublas become aliases of hip::host / roc::hipblas, every .cu (and the ggml MMQ
sources, with GGML_USE_HIP) is compiled as HIP. The native kernels' --use_fast_math (an nvcc flag) is not passed.

Correctness fixes found on the way (each one changed results or hung, not only speed):
- __fmul_rn / __fadd_rn / __fsub_rn contracted into FMA under HIP; they no longer do (as on NVIDIA) - the
  elementwise and quantize_act parity tests pass bitwise;
- the device check of the gfx906 build reads the card's gcnArchName (the wave32 check is the RDNA backend's);
- ggml MMQ got NVIDIA compute capability 9.0 for gfx906 and hung the prompt path; it gets the AMD arch code;
- __byte_perm is a scratch-memory array in HIP: the IQ4 table lookup uses llama.cpp's v_perm_b32 sequence
  (agent-request decode 29.9 -> 39.2 tok/s on 2x MI50);
- the stage profiler's wall clock runs at 25 MHz on gfx906, not 100 MHz;
- a CUDA toolkit-mismatch check that reads CUDART_VERSION is skipped (no CUDART_VERSION under HIP);
- gdn_rec_parity (cp.async, CUDA occupancy API) is not built for gfx906.

wave64 kernels (gfx906 only; the CUDA layout stays one env switch away for A/B, every one bitwise-checked):
- grouped native experts, STRATA_EXP_MODE=7 (default): gfx906 has no packed byte SIMD, so __vcmpne4/__vsub4
  became ~12 word ops each; signs are now dp4a(g ^ m, u) - dp4a(m, u) (exact in integers), grids and the group's
  activations sit in LDS, one weight load serves every token of the group. Same k-to-lane order, so the output is
  bitwise the old layout's (native_expert_bench, real GGUF rows). Verify window 60.0 -> 51.3 ms.
  STRATA_EXP_MODE=2 is the previous layout;
- MMVQ: one wavefront per row with a 64-lane butterfly, no LDS/barrier, same kernel for 1..8 columns
  (+2.3% decode; STRATA_MMVQ_WAVE=0: CUDA layout);
- router top-10 in one wavefront (register argmax, no block barriers), bit-identical (router_top10_parity);
- gr_down_multi split along K (8 rows x 5 slices of 2048 per block, fixed-order finish) (+2.9%;
  STRATA_GR_SPLIT=0: old);
- hyper-connection read: norm with its loads in flight and xn in registers, up with 8 lanes per row
  (51.3 -> 50.2 ms per window; STRATA_GR_FAST=0: old);
- the QSA prompt-attention tensor-core/WMMA kernels are not used on gfx906 (the FP32 kernel runs).

Benches (not ctest; they need a model shard): native_expert_bench (one expert layout against another, bitwise,
with timings) and fused_gr_bench (the fast hyper-connection kernels against the old ones).

2026-10-03 23:33:00 +02:00

..

hip_compat

hip: AMD gfx906 (Instinct MI50 / MI60, Radeon VII; wave64) as an opt-in build, STRATA_HIP_GFX906=ON

2026-10-03 23:33:00 +02:00

direct_file.hpp

C-4: the prompt's embedding in one gather, and the n-gram table read by several threads

2026-09-29 07:23:09 +02:00

memory.hpp

File tier: decide unbuffered reads from the expert bytes it reads, again once the RAM copy is built (#577)

2026-10-03 18:12:53 +02:00