[Vop3p]

On gfx103x the prompt path's dense GEMMs run through rocBLAS in FP16 (#835). rocBLAS picks the kernel for each 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. The GDN input projection (N = 10240, K = 2560) below T = 1152 tokens runs at 5.2 TFLOPS (8.1 ms at n = 799 on an RX 6900 XT); at T = 1152 rocBLAS's choice flips to a 38-TFLOPS kernel, and the best of the 237 solutions rocblas_gemm_ex_get_solutions lists for n = 799 runs it in 1.36 ms (6x). It is not an alignment effect: n = 800 is as slow as 799. In the engine's phase timing this is why a 799-token prompt's "gdn" phase (445 ms) was a third of an 8,874-token prompt's (1,345 ms).

This PR does for rocBLAS what tools/hip/tune_hipblaslt and STRATA_HIPBLASLT_TUNING do for hipBLASLt (which has no gfx1030 kernels):

• tools/hip/tune_rocblas.cpp: for each of the engine's 14 FP16 dense shapes at T = 256 ... 8192 (--case, --tokens change them), times rocBLAS's default kernel and every solution the library lists, checks the fastest against the default kernel's output (relative L2 within 2e-3, no element beyond 4 FP16 ulps of the largest output, nothing written outside the N x T result) and that they also run at T-1, T-37, T-63 and T/2+1 (a kernel can refuse a token count it was not measured at), and writes one row per shape and bucket: the winning solution index where it beats the default by --min-gain (5%), default otherwise. The default rows keep a neighbouring bucket's solution from reaching token counts where rocBLAS's own choice measured best.
• src/prefill/rocblas_tuning.hpp, Gemm::f16_inplace (gemm.cu): STRATA_ROCBLAS_TUNING names a table file, or a directory from which the engine takes <arch>-rocblas-<version>.txt for its own rocBLAS build. The table header carries the architecture and the full rocblas_get_version_string (tweak hash included) and any other pair is refused. The shape's nearest bucket decides; a tuned row's solution runs through a rocBLAS handle on the prompt stream, a default row or no row runs the kernel as before. One tiny default product at startup opens rocBLAS's kernel library (a solution-index call before that fails with internal_error although the index is valid - measured). A solution rocBLAS refuses at run time writes nothing and the default kernel runs that call; three refusals drop it for that shape. STRATA_ROCBLAS_VERBOSE=1 prints the pick per shape and a tuned_launches=... fallbacks=... summary.
• tools/hip/gfx1030-rocblas-5.6.0.8d1ae90e.txt: the table for ROCm 10.0.0's rocBLAS (/opt/rocm/lib), 112 rows, 76 with a solution. Note that a HIP build on Ubuntu 26.04 links Ubuntu's own librocblas5 (a 5.1 build) unless lib_dirs puts /opt/rocm/lib first, as setup's configuration does; 5.1's default kernel for the shape above is as slow, and its indices differ, so it would need its own table (tune_rocblas, about five minutes).
• setup.py: points the engine at tools/hip when a <arch>-rocblas-*.txt exists for the card; the engine's log then says rocBLAS tuning enabled (112 rows, 76 with a solution, gfx1030, rocBLAS 5.6.0.8d1ae90e) or why not.
• CMake: find_package(rocblas) (quiet) sets STRATA_ROCBLAS_AVAILABLE; tune_rocblas and the test build with it.
• tests/hip/prefill_rocblas_gemm.cpp (hip_prefill_rocblas_gemm, skips without STRATA_ROCBLAS_TUNING): refuses a table for another architecture or build, and checks Gemm's FP16 route against hipBLASEx (FP32 out) on the table's first two shapes at their bucket and at a smaller odd T. With STRATA_ROCBLAS_VERBOSE=1 it reported tuned_launches=4 fallbacks=0.
• docs/AMD_HIP.md: the RDNA2 section's new entry, a "rocBLAS table (gfx103x)" subsection under "Tuning table", a line in the setup section.

Measured (2x RX 6900 XT 16 GB, one card, Ryzen 5 5600X, 128 GB, Ubuntu 26.04, ROCm 10.0.0; 6f32ec0 + #835 + this, -DSTRATA_PREFILL_MMQ=ON; GSQ-RCO IQ3_S, --kv int8 --kv-resident 32768 --max-context 131072 --adapt-every 100000; STRATA_PREFILL_TIMING=1, temperature 0, max_tokens 8; one server per arm, four prompts from four offsets of one text; details and per-run lines in bench/results/2026-10-05-rdna2-rocblas-table/):

prompt tokens read	GPU timeline, default kernels	with the table	gdn phase	hc read phase
799	3,329 ms	2,985 ms (-10.3%)	443 -> 177	270 -> 77
495	2,477 ms	2,139 ms (-13.6%)	365 -> 146	81 -> 52
1,108	3,391 ms	3,395 ms (0%)	598 -> 201	158 -> 99
8,874	11,978 ms	11,459 ms (-4.3%)	1,346 -> 1,128	1,200 -> 898

The answers are the same in all four pairs. The 1,108-token prompt did not change because its saving reappeared as wait copy: that prompt is bound by the experts' PCIe transfer, and a faster GEMM only brings the GPU to the wait sooner. Prompts above ~1,150 tokens gain through their tail chunk and hc read only. Decode does not use these GEMMs.

ctest (-DSTRATA_BUILD_TESTS=ON, gfx1030, ROCm 10.0, STRATA_ROCBLAS_TUNING=tools/hip): 66 of 70 pass. hip_prefill_rocblas_gemm passes. Failed, for reasons outside this change: ple_parity (a Q2_0 PLE file that is not on this machine), expert_multi_test (this Zen 3 CPU has no AVX-512), platform_memory_test (mlock at the shell's default ulimit -l), expert_cache_segmented_test ("--vram-elastic (a segmented expert cache) is CUDA-only for now").

Builds on #835 (its commits are in this branch until #835 lands): the table serves the FP16-out route that #835 adds. One machine, one card, one model, one run per arm; another rocBLAS build needs its own table; gfx1031 / gfx1032 were not checked and the engine's architecture check refuses this table on them.

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.

[Niko1221]

Needs the AMD gfx103x tests and an A/B (10+ pairs, medians) before merging; the engine PRs are reviewed as a group. Queued, no GPU change in 0.1.40 without the numbers.

[Niko1221]

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

[Niko1221]

Thank you, xjc10. The rocBLAS solution table is deferred: 0.1.40 reworked #835 (it is opt-in), and #1006 now says #835 is faster on the RX 6800 than the table path. Rebase on 0.1.40, ten interleaved pairs, and an undrafted PR would bring it back. The PR was closed by the history rewrite.