Gyro-L (agentionai's 50 GiB file) stores its experts as a pair trellis: 128
values per block, a fp16 scale and 64 steps of K bits (K = 5, 6, 7), two
weights per step taken from the first pair of the TQ2_T codebook entry the
step's state selects. The engine refused the file because it did not know the
three type ids. This adds them next to the existing trellis types 144-146
without changing those.
What changed:
- gguf_reader.hpp / tools/gguf_reader.py: names and block sizes (42 / 50 / 58 bytes).
- tq_cpu.*, tq_states.hpp, tq_avx2.cpp: tq_supported / tq_kbits / tq_block_bytes
for the new ids (the bytes per block are 2 + 8K for the pair layout, 2 + 4K
for the others), a scalar pair kernel (the reference) and an AVX2 pair kernel
that is bit for bit equal to it; AVX-512 uses the AVX2 pair kernel.
- iq_kernels.cu: TqInfo, the lane decode (tqp_states8, a byte window), the
dot product, the decode-once Split, the fp16 dequantizer and the format lists
for 147-149. They are reached only by a pack that holds these types.
- moe_mmq.cu, strata_ggml_tq_overlay.cmake, third_party/ggml-tq/tq.cuh,
ggml-common-tq.h: the MMQ prompt path for the three types (overlay
enumerators, traits, tile loader branch, instances).
- native_expert_parity: names tqp5-7 and skips its synthetic check, with a
message, on a ggml that has no trellis traits.
- docs/GYRO.md: the "not supported" sentence now says how Gyro-L is read.
Changes to the contributor's patch:
- tq_cpu.cpp called the AVX2 pair kernel without the x86 guard the rest of the
file has, which breaks the Metal build (it compiles tq_cpu.cpp without
tq_avx2.cpp): the pair rows go through one helper that falls back to the scalar kernel.
- moe_mmq.cu named GGML_TYPE_TQP5-7 under STRATA_MMQ_TQ, but the public
agentionai llama.cpp has the older trellis types and not these: a new
STRATA_MMQ_TQP (set for the overlay and for a tree that has the tqp instances)
guards them, so that tree still builds. The HIP overlay lists get the three
instances too, and ggml_tq_overlay_test checks them (and sets the CMake
policy its IN_LIST needs).
- native_expert_parity: the skip test compared a type-name pointer to null, but
ggml_type_name of an id past the table reads past it (the test crashed on
an upstream ggml); it compares against GGML_TYPE_COUNT.
- docs/GYRO.md keeps the plain description and leaves out the contributor's
measured speeds (we have not repeated them).
Tests added: src/kernels/cpu/tqp_cpu_test.cpp (no ggml, no model: the scalar
kernel against a bit-by-bit reference from the format definition, every ISA
equal to scalar for 1-8 tokens and for gate/up rows, the type predicates for ids
0-299 against the old formulas), the pair types in tq_pascal_bench (fp16
dequantizer bit for bit against the definition, decode kernels within 3e-2,
decode-once == per-column), tools/test_gguf_tqp.py, and native_expert_parity
tqp5/6/7 registered when the ggml tree has the tqp instances.
Checked (no model): tqp_cpu_test passes on x86 (AVX2 on z440; AVX2 and
AVX-512 on the owner PC with MSVC); tq_pascal_bench on an RTX 5070 (sm_120,
MSVC build): 0 failures, the TQP dequantizer matches the definition in all
bits, and the checksum lines of TQ2_T/TQK6/TQK7 are byte-identical to the
build without this change. Builds: CUDA sm_86 (gcc), sm_60 (CUDA 12.9),
sm_120 (MSVC), SYCL, and the CPU files for aarch64.
Not covered: the contributor's own parity runs need a ggml tree that carries the
types (not public here), and the 50 GiB model run (load check, coherent
answer, 32K prompt, MTP, Gyro-S/M identity) is still to be done. Metal, SYCL and
HIP have no kernels for the types; HIP was read, not built.
A cmake -P script has no cmake_minimum_required, so CMake 3.x used the OLD behaviour and stopped at 'if given arguments ... IN_LIST: Unknown
arguments' (ctest ggml_tq_overlay with CMake 3.31). Checked with pip cmake 3.31.6 and 4.4.3 against the pinned llama.cpp: both print
'ggml_tq_overlay_test: ok'. A static test keeps the policy before the first IN_LIST in every cmake/*_test.cmake.
cmake/strata_ggml_tq_overlay.cmake writes, at configure time, a copy of the pinned ggml-cuda headers (and ggml.h, quantize.cu) into
the build folder with TQ2_T / TQK6 / TQK7 routed the way IQ2_XS is: same Q8_0_16 tile layout, dot kernels and tile configs (the
fork's own configs for them are copies of IQ2_XS's), three template instances, the tile loader and tq.cuh from third_party/ggml-tq.
Only the strata_mmq target sees the copy; every other type is compiled from the untouched files. Anchors are checked: another
llama.cpp turns the overlay off with a message (FP16 prompt path). moe_mmq.cu knows the trellis block geometry itself.
The tile loader and the trellis tables are agentionai/llama.cpp's (MIT, ggml authors; Laurent Zuijdwijk), ported from
agentionai/llama.cpp 93a24e5 ggml/src/ggml-cuda/{tq.cuh,mmq-load-tiles.cuh,tq-lut-i8.cuh}.
Checked: Gyro-S and Gyro-M exact_pp 4096/20000 residual dumps IDENTICAL to rc2 (R9700 gfx1201); Q2_0 and IQ3_XXS IDENTICAL to f-split
on gfx1201 and sm_86; ctest ggml_tq_overlay.
Three changes for two-card AMD PCs whose faster card does not come first in the
runtime's enumeration (Windows HIP orders the cards itself; HIP_VISIBLE_DEVICES
only filters):
- gfx1032 (RX 6600) joins the unvalidated build list - it shares the gfx1030 /
gfx1031 / gfx1034 dp4a path.
- STRATA_PRIMARY_DEVICE=N picks the visible ordinal the engine's own session
runs on; it is set before any context or allocation exists. The helper
caches take the other visible devices, the HIP arch check and the peer tier's
device queries follow it, and the primary's session is created on the chosen
device. A layer split still assumes device 0 and is refused with this set.
Unset or 0 is upstream's behavior unchanged.
- The helper-GPU caches are allowed beside the resident RAM mode: the experts a
helper holds are left out of the RAM copy as a split stage's are, and the two
adaptive tiers already keep each other's experts out. RemoteExperts::finish
spins on its stream like PeerExperts::finish - on HIP the CUDA branch's
cudaInitDevice(ScheduleSpin) does not exist.
Measured on an RX 6800 (16 GB, gfx1030) + RX 6600 (8 GB, gfx1032), Windows 11,
32 GB of RAM, Swift 1.5 IQ3_XXS at 160K context, the resident RAM mode with the
helper's cache left out of the copy: decode 37.5 -> 44.2 tok/s on 8K-token
requests with a 3800-slot helper; the CPU pool's share of the routed experts
falls from ~2.9 to ~0.4 entries per layer-window.
Rebased onto 296b59d2 (conflicts in generate.cpp, peer_experts.cpp, hip_backend.cmake, docs/AMD_HIP.md resolved by keeping main's code and adding the PR's). Changes: remote_caches is kept (still used by the --batch-mtp check); the spin in RemoteExperts::finish is HIP only (CUDA keeps its cudaInitDevice(ScheduleSpin) + blocking sync). NOT COMPILED YET, not run.
Co-authored-by: tuandat3019 <[email protected]>
rocBLAS picks the kernel for each FP16 product from its own table by shape, and on gfx1030 that pick is far
from the best kernel the same library holds for some shapes: N = 10240 K = 2560 (the GDN input projection)
below T = 1152 runs at 5.2 TFLOPS where the best of rocBLAS's 237 solutions for n = 799 is 6x faster (not an
alignment effect: n = 800 is as slow as 799).
tools/hip/tune_rocblas enumerates the library's solutions per shape and token bucket, checks each against the
default kernel's output and that it also runs at smaller odd token counts, and writes the winners (and
`default` rows bounding their reach) as tools/hip/<arch>-rocblas-<version>.txt; Gemm::f16_inplace runs the
table's solution for a shape (STRATA_ROCBLAS_TUNING: a file, or the directory from which the engine takes the
file of its own rocBLAS build; the version string, tweak hash included, is checked). One tiny default product
at startup opens rocBLAS's kernel library (a solution-index call before that fails although the index is
valid); a refused solution falls back to the default kernel for that call, after three refusals for good.
setup points the engine at tools/hip when a table for the card's architecture exists. Shipped: the gfx1030
table for ROCm 10.0.0's rocBLAS 5.6.0.8d1ae90e (112 rows, 76 with a solution).
Measured (RX 6900 XT, 6f32ec0 + #835 + this, temperature 0, the same answers with and without): a 799-token
prompt's GPU time 3,329 -> 2,985 ms (-10.3%), 495 tokens 2,477 -> 2,139 (-13.6%), 8,874 tokens 11,978 ->
11,459 (-4.3%); a PCIe-bound 1,108-token prompt unchanged (its saving reappeared as `wait copy`). ctest 66 of
70 (the 4 failures are outside this change); hip_prefill_rocblas_gemm, the new smoke test, passes with 0
fallbacks. Builds on #835. bench/results/2026-10-05-rdna2-rocblas-table has the per-run lines and the tuner's
summary.
(cherry picked from commit e35e103604)
Rebased onto 296b59d2: conflicts in setup.py (SETUP_ENV gained STRATA_ROCBLAS_TUNING beside the 0.1.44 entries) and docs/AMD_HIP.md (kept both sections) resolved.
Co-authored-by: xjc10 <[email protected]>
gfx1103 was missing from the gfx11 matrix-core lists, so the 780M ran the portable verify-window and prompt kernels. Added to the
compile-time lists and gfx_arch_is_gfx11_wmma(). Radeon 780M, Q2_0, 5 interleaved pairs: decode (code) 11.96 -> 18.25 tok/s,
same token ids; prompts unchanged. hip_prefill_wmma_gemm_parity passes on gfx1103 (3520 cases). Identity gate (defaults off,
mg_norepeat 10/10 against the portable build, exact_pp 4K/20K residual dumps identical).
The build stopped with a FATAL_ERROR because gfx1150 was in none of the architecture lists, although the device code
(the WMMA guards, sudot4) already covers it. It is now in the unvalidated list: the build warns and goes on. The
report's machine (96 GB, TheRock 7.14.1, IQ3_XXS: decode 17-19 tok/s, prompt about 125-135 tok/s) is in docs/AMD_HIP.md.
Setup's own AMD path is unchanged.
Co-authored-by: zeriyoshi <[email protected]>
gfx1010 (Navi 10) is measured on an RX 5700 XT: 8 GB card, Coder IQ1_M, 32K
context, prompt 115 / 144 / 149 tok/s and decode 18.0 / 20.3 / 23.3 tok/s at 4K /
16K / 30K prompt tokens, 6 of 6 needle checks at 8K and 30K, 61 of 61 ctest on the
card. Nothing beyond the arch list is needed: the portable scalar dp4a that main
falls to measured 162.7 / 24.3 against 158.3 / 24.0 with the __gfx1012__ SDWA
control (#442) - inside the run-to-run spread - and 16 of 16 HIP tests pass on it.
gfx1011 (Navi 12) is untested on hardware: same RDNA1 ISA, builds with 0 errors on
ROCm 7.14's gfx101X-dgpu wheels.
Two box facts, for whoever picks this up: hipBLASLt 1.4.0 refuses all 26 of the
engine's GEMM shapes on gfx1010 (heuristic_count = 0), so there is no tuning table
and prompts run on plain hipBLAS; and in a PC that also has an AMD iGPU, ROCm's HSA
runtime segfaults in GpuAgent::InitDma while enumerating the pair - below HIP, so
HIP_VISIBLE_DEVICES cannot reach it - and the engine must start with
ROCR_VISIBLE_DEVICES naming the one KFD node (rocminfo for the index).
Co-authored-by: AcckiyGerman <[email protected]>
CMake caches CMAKE_HIP_FLAGS_<CONFIG> while it enables HIP, from the
values its platform files set once the compiler is identified. When
HIP cannot compile CMake's identification source, that configure
caches them empty and then stops; every later configure of the same
build folder keeps the cached values, so the build that finally works
compiles every HIP source without -O, i.e. at -O0, while C and C++ get
-O3 -DNDEBUG. setup configures build-hip every time, so one failed
attempt is enough.
That happened on an RX 6800 (Windows 11, HIP SDK 7.2, CMake 4.0.1):
the first attempt failed on Visual Studio 2026's <cmath>, the next one
passed with the include fix, and the engine (0.1.30 with the Windows
build of #247, --expert-cache 2048) decoded at 0.24-0.27 tok/s with
7.1-7.5 s per decode window. STRATA_VERIFY_PROFILE put the time in
the GPU's own stages (per window: q8+qkv GEMVs 1468 ms and out-proj
849 ms in the DeltaNet layers, the LM head 540 ms); the waits for flags
the CPU raises were 41 ms of its 8.4 s per window. With the flags back:
59-62 ms per window, 28.1-28.3 tok/s.
After enable_language(HIP), hip_backend.cmake gives each empty
per-configuration HIP flag the value CMake computed for the compiler
(CMAKE_HIP_FLAGS_<CONFIG>_INIT), which is what a fresh build folder
gets. Flags with a value, from CMake or the command line, are left
alone; one set empty on purpose is filled too.
Reproduced on Windows (HIP SDK 7.2, CMake 4.0.1) with a first
configure that failed on --rocm-path flags split at the space in
C:/Program Files, and on Linux (CMake 3.28, TheRock's ROCm 7.10 clang)
with one that failed on a --rocm-device-lib-path that does not exist.
After a working second configure in the same folder,
CMAKE_HIP_FLAGS_RELEASE stays empty and native_mmvq.cu compiles
without -O; with this change it gets -O3 -DNDEBUG. A fresh folder gets
-O3 -DNDEBUG either way.
setup.py's engine fingerprint now covers cmake/ as well, so a change
there alone (this one) recompiles a locally built engine on its next
start; that includes NVIDIA engines compiled here.
From PR #313 (1e42cf3): src/prefill/wmma_gemm.cu/.h, the Gemm::f16/bf16 hooks before hipBLASLt, the
build macro STRATA_WMMA_GFX11 and its parity test; not its prompt attention (STRATA_HIP_WMMA covers it).
Fix: the intrinsics guard named only gfx1100/1101/1102 while the runtime gate takes any gfx11, so on
gfx1150/gfx1151 mma() compiled to "return c" (silently wrong); it now names gfx1150/gfx1151 too.
hip_prefill_wmma_gemm_parity on gfx1151: 3520 cases pass (f16 max_abs 2.4e-7, bf16 0).
CMake refuses to mix cl.exe with Clang for HIP, so on Windows ROCm's clang++
compiles the host code too (targeting the MSVC ABI, with MSVC's headers and
linker from vcvars). The CUDA->HIP shim is force-included with /FI when the
host front end is MSVC-style (cl / clang-cl) and with -include otherwise, its
path in forward slashes; CMAKE_HIP_ARCHITECTURES may also be space-separated.
On Windows the shim keeps <windows.h>'s min/max macros out (NOMINMAX,
WIN32_LEAN_AND_MEAN).
Taken from jagsan-cyber's #247 as ported by dvasdekis in #325 (2373c75); the
arch lists and the arch check stay 0.1.32's. Built on this PC with ROCm
10.2.0a20260930 (TheRock wheels, the version #325 ran on an RX 9070 XT) for
gfx1100;gfx1101;gfx1102;gfx1200;gfx1201;gfx1030: strata.exe and all 253
targets with the tests. CUDA builds do not read these files.
Co-authored-by: Dimitri Vasdekis <[email protected]>
dp4a through the plain signed v_dot4_i32_i8 (RDNA2 has no sudot4).
setup knows the card and the gfx103X-all TheRock index.
docs carry a community report (RX 6900 XT 16 GB: complete build, 28/28
ctest, 38-42 tok/s decode at 131,072 context, 330-339 tok/s prefill)
Cut down from #254 and #256 onto 0.1.30's AMD tables:
- setup.py: AMD_ARCHS / AMD_NAMES / ROCM_INDEXES / AMD_CARDS list gfx1101 (TheRock gfx110X-dgpu) and gfx1200
(gfx120X-all); amd_gpus() takes the sysfs root so the detection can be tested.
- cmake/hip_backend.cmake: gfx1101 and gfx1200 build without the "not validated" warning (a status line names
them community-validated); gfx1102 (#192: ctest only) keeps the warning.
- docs/AMD_HIP.md (Community-validated cards: the owners' setups, ctest and numbers) and README.
- tools/test_setup_amd.py: detection on a mocked KFD topology (every family, the CPU node, an iGPU and gfx1102
listed as not supported, product-name fallbacks, the ROCm index per family).
Left out: the PRs' own cmake/device.cu arch checks (0.1.30's compiled_for covers them), the dp4a change (main has
it), math_constants.h, the bench JSONs, the launcher, and both hipBLASLt tables (neither PR measured the prompt speed
without its table).
Co-Authored-By: Joseph Hohertz <[email protected]>
Co-Authored-By: efeisot <[email protected]>
- CMake: CMAKE_HIP_ARCHITECTURES may list gfx1100 and gfx1201 (validated) and
gfx1101 / gfx1102 / gfx1200 (build with a warning: same wave32, 64 KiB LDS and
dot4, not run by us; #192, #176). The list reaches the code as STRATA_HIP_ARCHS.
- Runtime: the engine (every GPU a run uses) and strata-device compare the card's
gcnArchName (up to the ':' feature suffix) with the compiled archs and require
wave32, with a message naming the card, its arch and the build's archs, instead
of a later "invalid device function".
- __dp4a uses __builtin_amdgcn_sudot4 (v_dot4_i32_iu8) on all five archs; the
portable fallback stays for other compilers/targets.
- setup.py: AMD_ARCHS gfx1100 + gfx1201; TheRock's wheel index per GPU family
(gfx110X-dgpu, gfx120X-all; the pinned 7.10.0a20251120 is on both), only that
family's libraries on the path, reinstall when the family changes; a system
ROCm older than 7.0 is passed over for the wheels; the engine is rebuilt when
the chosen card's arch is not in its BUILD.json. The hipBLASLt table is
tools/hip/<arch>-hipblaslt-<version>.txt for the card's arch AND the installed
hipBLASLt (read from hipblaslt-version.h, 1.4.1 -> 100401), else none with a
message; before, the highest-named table of the arch was taken whatever the
installed version, which the engine then refused.
doplxyz's minimal patch from #178 (their commit 0c59582): CMake, the
runtime device check, the signed dot4 intrinsic and setup's AMD_ARCHS
accept gfx1201 beside gfx1100.
Niko1221