HIP: gfx1201 (Radeon AI PRO R9700) support, faster packed byte intrinsics, and a gfx1201 hipBLASLt table by ttio2tech · Pull Request #262 · Niko1221/Strata

ttio2tech · 2026-09-30T16:01:47Z

Adds support for gfx1201 (RDNA4, wave32) alongside gfx1100 and speeds up the HIP packed-byte paths.

The build, setup and device check accept gfx1201 wave32 beside gfx1100.
__dp4a used v_dot4 only on gfx1100; gfx1201 fell back to a per-byte loop in every decode GEMV.
HIP's __byte_perm indexes a byte array in scratch memory; the IQ4_XS/Q2_0 table lookups now use v_perm_b32 (IQ4_XS attn_qkv 229 -> 34 us, Q2_0 attn_q 95 -> 23 us on the R9700).
__vsub4, __vsubss4 and __vcmpne4 work on all four lanes at once instead of a per-lane loop (Q3_K 14-24%).
tests/hip/intrinsics checks __byte_perm and the packed byte ops on 1,024 random operands.
tools/hip/gfx1201-hipblaslt-100401.txt: dense prompt GEMM solutions measured on the R9700.

Q2_0 decode: 58 -> 104 tok/s on the R9700. The output of a greedy CLI run is bit-identical before and after the intrinsic changes.

… packed byte ops as single instructions on HIP (Q2_0 decode 58 -> 104 tok/s on the R9700); a gfx1201 hipBLASLt table - The build, setup and device check accept gfx1201 wave32 beside gfx1100. - __dp4a used v_dot4 only on gfx1100; gfx1201 fell back to a per-byte loop in every decode GEMV. - HIP's __byte_perm indexes a byte array in scratch memory; the IQ4_XS/Q2_0 table lookups now use v_perm_b32 (IQ4_XS attn_qkv 229 -> 34 us, Q2_0 attn_q 95 -> 23 us on the R9700). - __vsub4, __vsubss4 and __vcmpne4 work on all four lanes at once instead of a per-lane loop (Q3_K 14-24%). - tests/hip/intrinsics checks __byte_perm and the packed byte ops on 1,024 random operands. - tools/hip/gfx1201-hipblaslt-100401.txt: dense prompt GEMM solutions measured on the R9700. The output of a greedy CLI run is bit-identical before and after the intrinsic changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…Cm 10.2.0a nightly (hipBLASLt 1.5.0) tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows, calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500). The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in Niko1221#262 (100401, hipBLASLt 1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both can be shipped side by side. Compared with Niko1221#262's 24 rows this table has the same 24 shapes (with solution ids chosen for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560 at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192). docs/AMD_HIP.md lists the shipped tables and how to check one.

…from #262 by ttio2tech) Cut down from #262 (ttio2tech): only the intrinsics and their test. - HIP's __byte_perm indexes a byte array in private memory (scratch traffic in the IQ4_XS/Q2_0 table lookups); byte_perm() is the same selection (nibble & 7 of the pair {y:x}) in one v_perm_b32. - __vsub4, __vsubss4 and __vcmpne4 work on all four byte lanes at once (Hacker's Delight 2-18) instead of a per-lane loop. - tests/hip/intrinsics checks them on 1,024 random operands against host references (passes on gfx1201). Left out: the arch/setup/device.cu parts (0.1.30 already builds and checks gfx1201; main's dp4a covers gfx1100-1102/1200/1201) and the gfx1201 hipBLASLt table (none is shipped for gfx1201). A/B vs main on amdbox (Coder IQ1_M, 128 greedy tokens): identity with a fixed cache (--expert-cache 1500 --prefill 2048), 5 + 5 starts per card, 4K / 16K / 4K again: the modal outputs are identical on the R9700 and on the RX 9070 XT (one start of 20 differed from its arm's mode, the known HIP start drift). Speed with setup's settings (auto cache/chunk), 2 + 2 starts per card, tok/s: | card | arm | 4K cold | 4K warm | 16K cold | 16K warm | decode | | R9700 | main | 1,027 | 1,792 | 1,695 | 1,812 | 46.0-52.2 | | | #262 | 1,023 | 1,786 | 1,681 | 1,808 | 53.0-61.0 | | RX 9070 XT | main | 781 | 1,018 | 1,421 | 1,517 | 32.7-41.4 | | | #262 | 780 | 1,017 | 1,424 | 1,517 | 34.9-44.0 | Decode +15% on the R9700, +5-7% on the 9070 XT; prompts unchanged. Co-Authored-By: ttio2tech <125389792+ttio2tech@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…plit GGUF, Q4_K/Q5_K/Q5_1/Q8_0 experts, RAM budget + SSD tier with routing prefetch), the low-RAM mode reads GGUF experts in place, Windows GGUF load 2x (#230), server status race (#266) and tokenizer (#268) fixes, truncated tool calls (#211/#231), AMD intrinsics (#262); setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

Niko1221 · 2026-10-01T05:14:33Z

Released in 0.1.31: the v_perm_b32 byte_perm and the SWAR __vsub4/__vsubss4/__vcmpne4 intrinsics with your test - decode +15% on the R9700 and +5-7% on the RX 9070 XT, identical output (5+5 starts per card). gfx1201 itself had shipped in 0.1.30, and we left out the hipBLASLt table (0-3% on gfx1201 here). Credited in the commit. Thank you!

The 0.1.30-era numbers predate the 0.1.31 decode work (the Niko1221#262 intrinsics and the hc-read rework). Same frozen methodology, fresh cells: decode +13-27% depending on tier, and the 32K prompt dip is gone (602 -> 1,565 tok/s, in line with its neighbours). Every cell traces to the sweep record; methodology unchanged: one-shot runs, greedy, 256 tokens, median of 3 non-stalled cells per tier (stalls stay recorded, excluded from the median).

…Cm 10.2.0a nightly (hipBLASLt 1.5.0) tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows, calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500). The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in #262 (100401, hipBLASLt 1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both can be shipped side by side. Compared with #262's 24 rows this table has the same 24 shapes (with solution ids chosen for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560 at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192). docs/AMD_HIP.md lists the shipped tables and how to check one.

Rebased onto current main: the README table's shape changed in 0.1.34 (the 128K-decode column was dropped), so the README row is now decode@4K 59 / prompt@32K 1,565 tokens/s. DETAILS keeps the full 1K-128K rows with the engine-version footnotes. Methodology unchanged: one-shot runs, greedy, 256 tokens, median of 3 non-stalled cells per tier (stalls recorded, excluded); every cell traces to the sweep record in our tree. Decode on 0.1.31 is +13-27% over the 0.1.30-era numbers (Niko1221#262's intrinsics), and the 32K prompt dip (602 on 0.1.30) is gone on 0.1.31 (1,565, in line with its neighbours).

README row adapted to the 0.1.34 two-column table shape (decode@4K 59 / prompt@32K 1,565 tokens/s); DETAILS keeps the full 1K-128K rows with the engine-version footnotes. Methodology unchanged: one-shot runs, greedy, 256 tokens, median of 3 non-stalled cells per tier. Decode is +13-27% over the 0.1.30-era numbers (Niko1221#262), and the 32K prompt dip is gone on 0.1.31.

…from #262 by ttio2tech) Cut down from #262 (ttio2tech): only the intrinsics and their test. - HIP's __byte_perm indexes a byte array in private memory (scratch traffic in the IQ4_XS/Q2_0 table lookups); byte_perm() is the same selection (nibble & 7 of the pair {y:x}) in one v_perm_b32. - __vsub4, __vsubss4 and __vcmpne4 work on all four byte lanes at once (Hacker's Delight 2-18) instead of a per-lane loop. - tests/hip/intrinsics checks them on 1,024 random operands against host references (passes on gfx1201). Left out: the arch/setup/device.cu parts (0.1.30 already builds and checks gfx1201; main's dp4a covers gfx1100-1102/1200/1201) and the gfx1201 hipBLASLt table (none is shipped for gfx1201). A/B vs main on amdbox (Coder IQ1_M, 128 greedy tokens): identity with a fixed cache (--expert-cache 1500 --prefill 2048), 5 + 5 starts per card, 4K / 16K / 4K again: the modal outputs are identical on the R9700 and on the RX 9070 XT (one start of 20 differed from its arm's mode, the known HIP start drift). Speed with setup's settings (auto cache/chunk), 2 + 2 starts per card, tok/s: | card | arm | 4K cold | 4K warm | 16K cold | 16K warm | decode | | R9700 | main | 1,027 | 1,792 | 1,695 | 1,812 | 46.0-52.2 | | | #262 | 1,023 | 1,786 | 1,681 | 1,808 | 53.0-61.0 | | RX 9070 XT | main | 781 | 1,018 | 1,421 | 1,517 | 32.7-41.4 | | | #262 | 780 | 1,017 | 1,424 | 1,517 | 34.9-44.0 | Decode +15% on the R9700, +5-7% on the 9070 XT; prompts unchanged. Co-Authored-By: ttio2tech <125389792+ttio2tech@users.noreply.github.com>

…plit GGUF, Q4_K/Q5_K/Q5_1/Q8_0 experts, RAM budget + SSD tier with routing prefetch), the low-RAM mode reads GGUF experts in place, Windows GGUF load 2x (#230), server status race (#266) and tokenizer (#268) fixes, truncated tool calls (#211/#231), AMD intrinsics (#262); setup requires it

…Cm 10.2.0a nightly (hipBLASLt 1.5.0) tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows, calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500). The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in #262 (100401, hipBLASLt 1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both can be shipped side by side. Compared with #262's 24 rows this table has the same 24 shapes (with solution ids chosen for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560 at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192). docs/AMD_HIP.md lists the shipped tables and how to check one.

…from Niko1221#262 by ttio2tech) Cut down from Niko1221#262 (ttio2tech): only the intrinsics and their test. - HIP's __byte_perm indexes a byte array in private memory (scratch traffic in the IQ4_XS/Q2_0 table lookups); byte_perm() is the same selection (nibble & 7 of the pair {y:x}) in one v_perm_b32. - __vsub4, __vsubss4 and __vcmpne4 work on all four byte lanes at once (Hacker's Delight 2-18) instead of a per-lane loop. - tests/hip/intrinsics checks them on 1,024 random operands against host references (passes on gfx1201). Left out: the arch/setup/device.cu parts (0.1.30 already builds and checks gfx1201; main's dp4a covers gfx1100-1102/1200/1201) and the gfx1201 hipBLASLt table (none is shipped for gfx1201). A/B vs main on amdbox (Coder IQ1_M, 128 greedy tokens): identity with a fixed cache (--expert-cache 1500 --prefill 2048), 5 + 5 starts per card, 4K / 16K / 4K again: the modal outputs are identical on the R9700 and on the RX 9070 XT (one start of 20 differed from its arm's mode, the known HIP start drift). Speed with setup's settings (auto cache/chunk), 2 + 2 starts per card, tok/s: | card | arm | 4K cold | 4K warm | 16K cold | 16K warm | decode | | R9700 | main | 1,027 | 1,792 | 1,695 | 1,812 | 46.0-52.2 | | | Niko1221#262 | 1,023 | 1,786 | 1,681 | 1,808 | 53.0-61.0 | | RX 9070 XT | main | 781 | 1,018 | 1,421 | 1,517 | 32.7-41.4 | | | Niko1221#262 | 780 | 1,017 | 1,424 | 1,517 | 34.9-44.0 | Decode +15% on the R9700, +5-7% on the 9070 XT; prompts unchanged. Co-Authored-By: ttio2tech <125389792+ttio2tech@users.noreply.github.com>

…plit GGUF, Q4_K/Q5_K/Q5_1/Q8_0 experts, RAM budget + SSD tier with routing prefetch), the low-RAM mode reads GGUF experts in place, Windows GGUF load 2x (Niko1221#230), server status race (Niko1221#266) and tokenizer (Niko1221#268) fixes, truncated tool calls (Niko1221#211/Niko1221#231), AMD intrinsics (Niko1221#262); setup requires it

…Cm 10.2.0a nightly (hipBLASLt 1.5.0) tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows, calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500). The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in Niko1221#262 (100401, hipBLASLt 1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both can be shipped side by side. Compared with Niko1221#262's 24 rows this table has the same 24 shapes (with solution ids chosen for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560 at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192). docs/AMD_HIP.md lists the shipped tables and how to check one.

…from Niko1221#262 by ttio2tech) Cut down from Niko1221#262 (ttio2tech): only the intrinsics and their test. - HIP's __byte_perm indexes a byte array in private memory (scratch traffic in the IQ4_XS/Q2_0 table lookups); byte_perm() is the same selection (nibble & 7 of the pair {y:x}) in one v_perm_b32. - __vsub4, __vsubss4 and __vcmpne4 work on all four byte lanes at once (Hacker's Delight 2-18) instead of a per-lane loop. - tests/hip/intrinsics checks them on 1,024 random operands against host references (passes on gfx1201). Left out: the arch/setup/device.cu parts (0.1.30 already builds and checks gfx1201; main's dp4a covers gfx1100-1102/1200/1201) and the gfx1201 hipBLASLt table (none is shipped for gfx1201). A/B vs main on amdbox (Coder IQ1_M, 128 greedy tokens): identity with a fixed cache (--expert-cache 1500 --prefill 2048), 5 + 5 starts per card, 4K / 16K / 4K again: the modal outputs are identical on the R9700 and on the RX 9070 XT (one start of 20 differed from its arm's mode, the known HIP start drift). Speed with setup's settings (auto cache/chunk), 2 + 2 starts per card, tok/s: | card | arm | 4K cold | 4K warm | 16K cold | 16K warm | decode | | R9700 | main | 1,027 | 1,792 | 1,695 | 1,812 | 46.0-52.2 | | | Niko1221#262 | 1,023 | 1,786 | 1,681 | 1,808 | 53.0-61.0 | | RX 9070 XT | main | 781 | 1,018 | 1,421 | 1,517 | 32.7-41.4 | | | Niko1221#262 | 780 | 1,017 | 1,424 | 1,517 | 34.9-44.0 | Decode +15% on the R9700, +5-7% on the 9070 XT; prompts unchanged. Co-Authored-By: ttio2tech <125389792+ttio2tech@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…plit GGUF, Q4_K/Q5_K/Q5_1/Q8_0 experts, RAM budget + SSD tier with routing prefetch), the low-RAM mode reads GGUF experts in place, Windows GGUF load 2x (Niko1221#230), server status race (Niko1221#266) and tokenizer (Niko1221#268) fixes, truncated tool calls (Niko1221#211/Niko1221#231), AMD intrinsics (Niko1221#262); setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…Cm 10.2.0a nightly (hipBLASLt 1.5.0) tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows, calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500). The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in Niko1221#262 (100401, hipBLASLt 1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both can be shipped side by side. Compared with Niko1221#262's 24 rows this table has the same 24 shapes (with solution ids chosen for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560 at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192). docs/AMD_HIP.md lists the shipped tables and how to check one.

Default off (STRATA_RETAINED_PROMPT=1). One tagged request and prompt submission across BQSUSPEND/BQRESUME/BQRESUMED; exact plan/grid/cursor retained, drained preflight, request-local stale/bad-cursor refusal, suspended cancellation and final-issued MTP catch-up. The frozen screen is negative: 7.599 vs 8.910 committed tok/s, decoder gap 2.069 vs 1.934 s; not promoted.

bsorensen110 mentioned this pull request Oct 1, 2026

hip: fast packed-byte intrinsics for RDNA3/RDNA4 (v_perm_b32 + SWAR) #322

Closed

bsorensen110 mentioned this pull request Oct 1, 2026

hip: hipBLASLt tuning table for gfx1201 (R9700) with ROCm 10.2.0a nightly (hipBLASLt 1.5.0) #339

Closed

Niko1221 closed this Oct 1, 2026

Niko1221 mentioned this pull request Oct 1, 2026

Docs: AMD RX 7900 XTX (gfx1100) IQ3_S rates row (1K-128K) #319

Closed

jmnargi mentioned this pull request Oct 1, 2026

Merge upstream v0.1.31 with V100 (Volta) support preserved jmnargi/Strata-V100#7

Merged

xyzzing mentioned this pull request Oct 1, 2026

HIP/ROCm (gfx1100) port: the full IQ3_S model builds, runs and serves on an RX 7900 XTX — findings and what we would like to send #106

Closed

xyzzing mentioned this pull request Oct 7, 2026

prompt attention: STRATA_PROMPT_ATTN_IMMA=1 - the int8-KV kernel on INT8 tensor cores (opt-in, closer to FP32) #1367

Closed

Navigation Menu

Repository navigation

HIP: gfx1201 (Radeon AI PRO R9700) support, faster packed byte intrinsics, and a gfx1201 hipBLASLt table - #262

HIP: gfx1201 (Radeon AI PRO R9700) support, faster packed byte intrinsics, and a gfx1201 hipBLASLt table#262
ttio2tech wants to merge 1 commit into
Niko1221:mainfrom
ttio2tech:hip-gfx1201

ttio2tech commented Sep 30, 2026

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

ttio2tech commented Sep 30, 2026

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants