Repository navigation
Conversation
2a33f5b to
1079741
Compare
|
Rebased on 0.1.40.3 (d5ea713); the conflicts were with the gfx1103 THP lines in setup.py and the hipBLASLt calibration notes in AMD_HIP.md (both kept). One small docs commit on top, 4cab840: the shipped gfx1030 table was measured on an RX 6900 XT. Other gfx1030 cards (RX 6800 / 6800 XT / 6950 XT, Radeon PRO W6800 / V620) can use it as it is, since the solution indices are valid on every gfx1030, and can make their own with An earlier version of this comment described a card-width field and a per-device cache. I dropped both to keep this PR to the table. #1151 no longer depends on this PR. Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030, PCIe 4.0 x8 each) / Ryzen 5 5600X, ROCm 10.0. |
fa356d9 to
4cab840
Compare
What the release brought for this box (Niko1221#1709 Windows HIP VRAM floor, the HIP stream-capture fix, the Linux-only PCIe telemetry), the merge table with what was kept and why Niko1221#1439/Niko1221#1841/Niko1221#1656 went, the measurements (0.1.42+local +5.6% at 16K; the four new PRs +4.4% on top), and the ranked shortlist for the next cycle (Niko1221#1297 bisect, Niko1221#1909, Niko1221#1324, Niko1221#1167, Niko1221#1298/Niko1221#1151) with the gfx1100 reading list.
…er + a gfx1030 table) Merged cleanly. The engine picks <arch>-rocblas-<version>.txt for its own card and rocBLAS build, or runs the default kernels and says so, so on gfx1100 without a table it is inert - its measured win is RDNA2.
…10.9% Six PRs merged individually (Niko1221#1297 Niko1221#1909 Niko1221#1324 Niko1221#1167 Niko1221#1298 Niko1221#1151), none of which changes the default path (the arms are opt-in and Niko1221#1297 needs a tuning table). Reading the issue tracker produced the real win: issue Niko1221#1549 documents ring_default_slots() sizing the staging ring at 384 (1024 for a fused ring) when g_pinned_share >= 0.9, and our STRATA_PREFILL_RING=96 override - left from tuning against a smaller budget - was costing 16.5% at 4K and 10.9% at 16K. The override is now gone, so the engine's own rule applies. Measured, with the LRU left off (it buys prefill but eats the resident budget: 11.9 -> 9.1 GiB).
…p-server.bat fixed A (in-tree levers): - Niko1221#1297 kernel tuning: RAN. The tuner kept the DEFAULT shape in all nine cases (3 expert shapes + 6 mmvq shapes) and wrote a table with 0 changes - nothing to gain on gfx1100. The tuner also needs its binary copied next to the engine on Windows (autotune.py looks for a literal 'tune_decode'). - Niko1221#1151 SIMT scorer: ADOPTED (STRATA_SELECT_SIMT=1) - +3.9% at 16K and +4.5% at 4K on top of the WMMA select, measured at matched budgets after the leak below was cleared. - Niko1221#1324 LRU: off. 2 GiB is a wash (16K -0.8%, budget -1.4 GiB); 6 GiB costs 14% and crushes the budget to 4.9 GiB. - Niko1221#1167 rocBLAS table: inert on gfx1100 - prompt_f16() is gfx103x-only ('on gfx103x it is ~2x faster'). - A5 hipBLASLt rows: deferred (our 45-row table already beats upstream's 34; the tuner needs a GPU-exclusive run). B: Niko1221#1838/Niko1221#1803/Niko1221#1876 applied and NEUTRAL (-2..+2% at both sizes, so kept); its conversational test shows no measurable difference either. C: Niko1221#1913 does not reproduce here - per-entry host cost is flat at 0.77-0.86 ms (their box jumps 3-5x). Niko1221#1903's placement axis answered: --expert-cache-per-layer costs 3%. Issue Niko1221#1910 does not apply: the interleaved 4K pairs put 0.1.42 +2.1% AHEAD of 0.1.41 (923.6/923.5/962.2 vs 958.7/957.4/953.2). Niko1221#1906 answered: reuse works, up to 96% of a turn's prompt reused (1286/1341). Bug found and fixed: stop-server.bat killed a fixed list of old binary names, so newer A/B engines leaked - one held 11.4 GiB and crushed the next measurement's RAM budget to 0.0 GiB. It now kills any strata*.exe (matching tasklist) plus tune_decode; that leak is what made the first group-B/SIMT/per-layer readings unusable.
…he GDN projection 7x below 1,152 tokens; short prompts -10 to -14%) 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 + Niko1221#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 Niko1221#835. bench/results/2026-10-05-rdna2-rocblas-table has the per-run lines and the tuner's summary. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> (cherry picked from commit e35e103)
…900 XT; other gfx1030 cards can make their own with tune_rocblas (reports welcome) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
4cab840 to
db64983
Compare
|
Rebased on 0.1.43.1 ( Ten interleaved pairs on 0.1.43.1 + this PR, one binary, one RX 6900 XT (gfx1030, ROCm 10.0 / rocBLAS 5.6.0), IQ3_S,
The ten pairs are within +-0.6 % of each other (1K: +13.7 to +14.1 %, 4K: +11.5 to +12.6 %). On #1006: the comparison there (RX 6800, "#835 is faster") is between #1006's FP32-SGEMM route and #835's FP16 route. This table is not a third route: it picks rocBLAS's solutions for #835's own FP16 GEMMs, so it only runs when Other architectures: a gfx1100 fork took this PR (Maddoxkkm/Strata@5c521a6c) and reports it inert there - the engine only loads ctest (gfx1030, this branch, ROCm 10.0's runtime via Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030, PCIe 4.0 x8 each) / Ryzen 5 5600X, ROCm 10.0. |
|
Thanks Vop3p. Ported with your commits (0.1.45, opt-in). We have no gfx103x card, so the +12-14% is your data; what we proved is that it is inert elsewhere: base and the new build give identical answers on an R9700 (10 prompts plus exact 4K/20K) and the same speed (decode 1.001, prompt 1.002 on Q2_0). It stays opt-in; if a second RX 6000 owner reports the same gain we would look at making it the default for gfx1030. |
|
Thanks for the pull request. Your change is already in If part of your work did not make it in and you want it considered, open a new pull request on top of |
2 participants
Re-filed from #981 (closed with the history rewrite) as the one commit on top of #835, which is in 0.1.40.1. The table serves the
STRATA_HIP_PROMPT_F16=1route; with it off nothing runs through it. Inf16_inplacethe table's kernel is tried first,cublasGemmExotherwise, and the review'sSTRATA_DBG_NANcount follows either. The tool (tools/hip/tune_rocblas), the table format and the engine route are described in #981; the table file is the one measured there (tools/hip/gfx1030-rocblas-5.6.0.8d1ae90e.txt).Measured on 0.1.40.1 (
82f46a8+ this commit; RX 6900 XT, one card, ROCm 10.0.0 / rocBLAS 5.6.0.8d1ae90e;STRATA_HIP_PROMPT_F16=1 STRATA_SH_STREAM=1,--kv int8 --kv-resident 32768 --max-context 131072 --spec 4; the #884 GFXOFF workaround in place, which does not change speed):hip_prefill_rocblas_gemm(ctest,STRATA_ROCBLAS_TUNING=tools/hip): PASS -rocBLAS tuning enabled (112 rows, 76 with a solution, gfx1030, rocBLAS 5.6.0.8d1ae90e), 4 cases within FP16 rounding of hipBLASEx,tuned_launches=4 fallbacks=0. On 0.1.40.1 the test needsSTRATA_HIP_PROMPT_F16=1as well, since the FP16 route is opt-in now; without it the test skips.max_tokens8, temperature 0: 12 pairs per length, the same slice in each pair. Prompt read speed from the engine's own line:--short-read 640the prompt is read token by token, no prompt GEMMsDecode is not touched (the table serves the prompt GEMMs only). The engine logs
rocBLAS tuning enabled (...)when the table is read, so an arm's state is visible in its log.Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030, PCIe 4.0 x8 each) / Ryzen 5 5600X, ROCm 10.0.