Repository navigation
Conversation
Tuned a8w8 bpreshuffle tables for the Kimi-K3 decode and dense FP8 linears and a tuned bf16 table for its decode GEMMs (TP8, gfx950), plus the untuned shape lists they were tuned from (dense FP8, prefill bf16). The new decode bf16 table retunes the M=8/16/32/48/64, N=896, K=3584 shapes, so their rows are removed from kimik3_bf16_tuned_gemm.csv to keep one row per shape in the bf16_tuned_gemm merge. Co-Authored-By: Devin AI <devin-ai-integration[bot]@users.noreply.github.com> (cherry picked from commit 3be28ce)
github-actions
Bot
changed the title
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
One backend per PR: PR title tags & labels: |
|
Hi @aryatsc, heads-up: #6410 adds the same 20 gfx950 keys for N=6144, K=7168 (M=1, 2, 4, ..., 8192) in Whichever PR merges second should drop its copy of those 20 rows. One note: this PR's M=8192 row uses the FlyDSL 8-wave kernel, and because of padded-M lookup it also serves every M above 8192. #6410's rows avoid the 8-wave kernel and measure 1.41x vs main on this shape. Happy to coordinate either way. |
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Thanks for catching this. We dropped our copy of the 20 N=6144, K=7168 rows (6bd43d0), so #6410 owns that shape. test_config_shape_collision passes on this branch. With #6410's kimik3 file added, it reports 9 duplicates between that file and dsv3_a8w8_bpreshuffle_tuned_gemm.csv, none from this PR. |
| gfx,cu_num,M,N,K,bias,dtype,outdtype,scaleAB,bpreshuffle,libtype,solidx,splitK,us,kernelName,err_ratio,tflops,bw | ||
| gfx950,256,16,144,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,0,8,4.7084,flydsl_hgemm_abf16_wbf16_bf16_t16x16x128x8_ks8_w1x1x1_bias0_ktail0_gm0_pft_gfx950,0.009,7.02,488.14 | ||
| gfx950,256,24,144,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,0,8,5.0605,flydsl_hgemm_abf16_wbf16_bf16_t32x16x128x5_ks8_w2x1x1_bias0_ktail0_gm0_pft_gfx950,0.0082,9.79,477.29 | ||
| gfx950,256,32,144,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,0,8,5.0922,flydsl_hgemm_abf16_wbf16_bf16_t32x16x128x4_ks8_w2x1x1_bias0_ktail0_gm0_pft_gfx950,0.0076,12.97,497.3 |
There was a problem hiding this comment.
better keep only 1 tuned_gemm.csv . Agent can help extract certain shapes from it. Same as a8w8.
…ig files Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
WAITING ON AUTHORPR Self-ReviewRevision Waiting for the PR author. The readiness report will appear after the self-review answers are complete for this revision. Questions to answer in the PR description
Edit the Missing or invalid answers
Shadow mode reports what would block but keeps the status successful. Set |
Done in 5b96cdb. The decode/dense rows now live in the existing kimik3 files, one tuned and one untuned file per kind:
The separate kimik3decode/kimik3dense/decode/prefill/dense files are gone. test_config_shape_collision reports no duplicate shapes for these files. |
- Replace the 10 hipblaslt rows (M=8..64 N=7168 K=43008, M=40/56 N=20480 K=7168) with torch/flydsl winners - Retuned with gemm_a16w16_tune.py on gfx950; 0.9-5.7% slower than same-session hipblaslt - Verified all 10 rows through --run_config (status OK); test_config_shape_collision passes Co-authored-by: Yu-Yun Chang <yuyun.chang@amd.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
waqahmed-amd-fi
left a comment
waqahmed-amd-fi
left a comment
There was a problem hiding this comment.
Static review of bafbd99: nothing was run, and every timing below is a us value read from the CSVs. GitHub CI ran only the title-tag and self-review workflows on this head; the test workflows wait for maintainer approval because the branch is on a fork. The description quoted is the version edited on Oct 29 at 21:11 UTC.
Tasks 1 and 2 should be resolved before merge. Task 3 is at your discretion, and the design question does not block.
Tasks
- Drop the 11 N=6144, K=7168 rows, or A/B them against the rows from #6410
- Correct the description's claims about which selections change
- A/B the switch to Triton for M=17–32 and M=40 at N=1536, K=128
Lower importance
- Nit: Keep the untuned lists in step with the tuned rows
Design questions
- Were the a8w8 rows tuned without
--splitK? Neither approach is presumed wrong.
Heads-up
- Seven
hipblasltrows (N=6016, K=7168, lines 52, 54, 56 and 65–68 ofkimik3_bf16_tuned_gemm.csv) predate this PR and have the persistence problembafbd99fixed for its own rows, so they could be a follow-up.
| gfx950,256,24,6016,7168,torch.float8_e4m3fn,flydsl,2375,0,10.9996,flydsl_bpreshuflle_32x32x1024_F8_F8_B16_0x0x0x1_default,188.18,3962.28,0.0 | ||
| gfx950,256,32,6016,7168,torch.float8_e4m3fn,flydsl,2315,0,11.7482,flydsl_bpreshuflle_32x32x1024_F8_F8_B16_0x0x0x2_default,234.92,3722.88,0.0 | ||
| gfx950,256,40,6016,7168,torch.float8_e4m3fn,cktile,37,0,14.9915,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x2_32x64x512_1x4x1_16x16x128_default,230.12,2927.7,0.0 | ||
| gfx950,256,48,6016,7168,torch.float8_e4m3fn,cktile,37,0,15.3141,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x2_32x64x512_1x4x1_16x16x128_default,270.32,2876.06,0.0 |
There was a problem hiding this comment.
Design question: Were the a8w8 rows tuned without --splitK?
All 53 new a8w8 rows have splitK 0 and no _ksN kernel, and the tuner adds FlyDSL split-K candidates only under --splitK, which is off by default (gemm_a8w8_bpreshuffle_tune.py:543, base_tuner.py:141). At the 20 N=6144, K=7168 keys this PR carried until 6bd43d0, its picks matched #6410's within 2.3% where both chose the same kernel, but were 1–4% slower at M=32, 128 and 256 and 7–23% slower at M=48, 72, 80 and 96, where #6410 picked split-K kernels. At M=48 the slower pick was the same 32x64x512 CKTile tile as this row, at 15.09 µs against 13.66 µs.
If --splitK was off, a re-tune with it may find faster kernels for these N=6016 rows and the N=768 ones. Those shapes had no tuned rows before, so this need not block, and leaving split-K out on purpose is also fine.
There was a problem hiding this comment.
Are deterministic solutions all you need?
| gfx950,256,4096,6016,7168,torch.float8_e4m3fn,flydsl,1000033,0,151.1194,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_4x0,2337.63,805.76,0.0 | ||
| gfx950,256,8192,6016,7168,torch.float8_e4m3fn,flydsl,1000028,0,244.624,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_2x4,2888.2,819.25,0.0 | ||
| gfx950,256,16384,6016,7168,torch.float8_e4m3fn,flydsl,1000031,0,481.1443,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_3x4,2936.84,743.43,0.0 | ||
| gfx950,256,40,6144,7168,torch.float8_e4m3fn,cktile,151,0,14.9501,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x3_32x64x512_1x4x1_16x16x128_default,235.67,2997.87,0.0 |
There was a problem hiding this comment.
Task 1: Drop the 11 N=6144, K=7168 rows, or A/B them against the rows from #6410 (category: performance)
Lines 1207–1217 add 11 rows for the shape the Oct 28 reply said #6410 owns. Because the lookup also tries padded M (gemm_op_a8w8.py:581), they take over M values that #6410's rows serve today: M=40 from its M=48 row (line 1350), 88 from M=96 (line 1355), 97–112 and 120 from M=128 (line 1356), and 177–192, 353–384, 737–768, 1473–1536, 3009–3072 and 16257–16384. At M=40, 88, 97–112 and 120, #6410's kernels need no more M tiles than at their own M and record 13.66, 16.13 and 18.70 µs; the new rows record 14.95, 16.80 and 20.52–20.55 µs. The two tuning runs agree within 2.3% wherever they chose the same N=6144 kernel (this PR's rows before 6bd43d0), and at M=48 this PR's earlier pick, the CKTile kernel this row uses, recorded 15.09 µs against #6410's 13.66 µs. So those M probably run 4–10% slower, though nothing here was measured. The collision guard passes because it compares exact keys.
Could these 11 rows be dropped? Keeping a row is fine if a back-to-back A/B of gemm_a8w8_bpreshuffle against main, alternating rounds as #6410 did, shows it faster at every M it takes over.
| gfx950,256,16384,7168,1792,False,torch.bfloat16,torch.bfloat16,False,False,torch,0,0,310.0445,native,0.0,1357.57,1029.82 | ||
| gfx950,256,32768,7168,1024,False,torch.bfloat16,torch.bfloat16,False,False,torch,0,0,382.655,native,0.0,1257.1,1441.38 | ||
| gfx950,256,32768,7168,1792,False,torch.bfloat16,torch.bfloat16,False,False,torch,0,0,622.2538,native,0.0,1352.85,984.96 | ||
| gfx950,256,8,896,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,403,7,5.4581,flydsl_hgemm_abf16_wbf16_bf16_t16x32x64x8_ks7_w1x1x1_bias0_ktail0_gm0_pft_gfx950,0.0197,9.41,1189.83 |
There was a problem hiding this comment.
Task 2: Correct the description's claims about which selections change (category: docs)
Answering the self-review question for config PRs, the description says these rows are consumed only by lookups "whose full key matches a row" and that "every other shape keeps its previous selection", citing the collision guard. Both lookups also try padded M (tuned_gemm.py:134, gemm_op_a8w8.py:581), so a row also serves the M values that pad to it, and the guard compares exact keys only. At this head 23 new rows take over M values from main's rows (Task 1 has the N=6144 ones), and the diff replaces 8 existing bf16 rows, starting with this deleted one; the description mentions neither. It also names no tuning stack or tuner flags for the 75 rows from 04db8eb; only the 10 retuned rows have a --run_config result.
Could you rewrite that answer from a per-M comparison of the selected row on main and on this head, list the replaced rows, and name the stack and tuner flags behind the 75 older rows?
| gfx950,256,24,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,triton,0,0,2.82,auto,0.0,3.35,167.76 | ||
| gfx950,256,32,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,triton,0,0,2.36,auto,0.0,5.33,211.74 | ||
| gfx950,256,40,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,triton,0,0,3.23,auto,0.0,4.87,162.95 |
There was a problem hiding this comment.
Task 3: A/B the switch to Triton for M=17–32 and M=40 at N=1536, K=128 (category: performance, at your discretion)
These rows move M=17–32 and M=40 at N=1536, K=128 from main's FlyDSL kernel (t16x16x64x3_ks1, whose M=32 and M=48 rows record 2.73 and 2.76 µs) to Triton, which records 2.82 µs at M=24, 2.36 at M=32 and 3.23 at M=40. That kernel needs no more 16-row tiles at M=24 and 40 than at 32 and 48, and in this tuning run the same FlyDSL kernels read 5–8% faster than main's rows did (line 749 against the deleted line 630, and line 762 at M=40 against the deleted M=48 line 697). So M=24 and M=40 probably get slower and M=32 may get faster; Triton's times here fall from M=24 to M=32, so they look noisy, and nothing here was measured.
Could you time Triton and main's FlyDSL kernel back to back at M=24, 32 and 40, as #6410 did for the DSv3 rows it replaced, and keep main's rows where Triton does not win? The same check would settle the other five kernel switches (lines 750, 752, 754, 756 and 763), whose 1.9–6.8% CSV gains are within that drift.
| 1024,12,7168,False,torch.bfloat16,torch.bfloat16,False,False | ||
| 2048,12,7168,False,torch.bfloat16,torch.bfloat16,False,False | ||
| 4096,12,7168,False,torch.bfloat16,torch.bfloat16,False,False | ||
| 16384,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False |
There was a problem hiding this comment.
Nit: Keep the untuned lists in step with the tuned rows. 14 of the 15 shapes added from here have no tuned row, and 35 of the new tuned rows (24 bf16, 11 a8w8) have no untuned row, so a --all re-tune would skip those 35 and tune the 14; #6410 added an untuned row for every tuned key.
- Remove the 11 a8w8 bpreshuffle rows at N=6144, K=7168 (and their untuned lines). With M padding they took over M values from ROCm#6410's rows. - Untuned lists now hold exactly the (M, N, K) of the tuned rows this PR adds. Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
|
Thanks, this was a thorough review. Changes are in 2147087:
Untuned lists: they now hold exactly the (M, N, K) of the tuned rows this PR adds.
|
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
|
Re your comment, checked at
Nothing here blocks merge from my side. |
waqahmed-amd-fi
left a comment
waqahmed-amd-fi
left a comment
There was a problem hiding this comment.
Review at 77eb4f9 is approved from my side, no blocking issues.
Left to the author: an optional A/B of the N=1536, K=128 Triton rows and three small description fixes.
|
Thanks for the re-check and approval. Description updated with your three fixes: 64 rows without a |
Motivation
Kimi-K3's decode and dense-FP8 linears on gfx950 (TP8) fall back to untuned a8w8/bf16 GEMM configs.
Modifications
kimik3_bf16_*anda8w8_bpreshuffle_*_kimik3tuned/untuned files. No new files.Tests
test_config_shape_collisionreports no duplicate shapes for these files. Used by the Kimi-K3 static FP8 dense PR in SGLang.PR Self-Review
Reviewed revision
2147087
Behavior and intent
Only tuned-config CSVs change: 42 a8w8 bpreshuffle rows and 24 bf16 rows (8 of the bf16 rows replace existing kimik3 rows at the same key). No code changes.
Validation and unverified scope
python3 -m unittest op_tests.tuning_tests.test_config_shape_collisionin the MI355X serving image: no duplicate-shape failures in the bf16 or a8w8 bpreshuffle families. (Two other tests in that file need neweraiter/jitcode than the image and fail the same way on cleanmain.)gemm_a16w16_tune.py --run_config(status OK, @yuychang).gemm_a8w8_bpreshuffle_tune.py,gemm_a16w16_tune.py) on our Kimi-K3 serving image. The exact tuner commit and flags for these rows were not recorded; every a8w8 row hassplitK=0. They have no--run_configreplay on the currentmain.PR-specific question
Lookups are gfx950 / cu_num 256 only, keyed on bf16
(M, N, K, bias, dtype, outdtype, scaleAB, bpreshuffle)and a8w8 bpreshuffle(M, N, K, q_dtype_w). A row also serves nearby M values:tuned_gemm.pyandgemm_op_a8w8.pytry exact M, thengetPaddedMlevel 0 (next multiple of 16/32/64/128), then level 1 (next power of two, capped at 8192 when N > 4096). So a new row changes the selection for every M that pads to it before reaching an existing row.The table below is the full set of selection changes. It replays that lookup for M = 1..32768 over the merged family (
configs/<family>.csvplus everymodel_configs/*<family>*.csv) onmainand on this head. Shapes not listed keep their selection.main(N=768/6016 K=7168 a8w8; N=144 K=7168 and N=7168 K=43008 bf16), so the change there is default → tuned.mainwould pick (recorded 3.23 µs vs FlyDSL M=48 row 2.76 µs, not an A/B).triton_gemmignores the row's config, so M=56 makes the same Triton call either way, and M=17-32 follow the M=24 Triton row. Every other change has a lower recorded time.Changed op tests
None
Per-M selection, main vs this head
a8w8_bpreshuffle N=768 K=7168
a8w8_bpreshuffle N=3584 K=7168
a8w8_bpreshuffle N=6016 K=7168
a8w8_bpreshuffle N=7168 K=512
a8w8_bpreshuffle N=7168 K=1792
bf16 N=144 K=7168
bf16 N=896 K=3584
bf16 N=1536 K=128
bf16 N=7168 K=1536
bf16 N=7168 K=43008
bf16 N=20480 K=7168