Skip to content

[Config] Add Kimi-K3 MI355X GEMM tuning tables #6959

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
aryatsc wants to merge 11 commits into
base: main
Choose a base branch
from
+151 −8
Open

[Config] Add Kimi-K3 MI355X GEMM tuning tables #6959

Show file tree
Hide file tree
Changes from 8 commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
04db8eb
Add Kimi-K3 MI355X GEMM tuning tables
aryatsc Oct 13, 2026
200ccc4
Merge branch 'main' into devin/k3-gemm-tables-1793030216
aryatsc Oct 26, 2026
78373be
Merge branch 'main' into devin/k3-gemm-tables-1793030216
aryatsc Oct 26, 2026
6bd43d0
Drop the N=6144, K=7168 rows that #6410 tunes
aryatsc Oct 28, 2026
14ff2cc
Merge branch 'main' into devin/k3-gemm-tables-1793030216
aryatsc Oct 28, 2026
5b96cdb
Fold Kimi-K3 decode and dense GEMM rows into the existing kimik3 conf…
aryatsc Oct 29, 2026
952fa0e
Merge branch 'main' into devin/k3-gemm-tables-1793030216
aryatsc Oct 29, 2026
bafbd99
fix: retune Kimi-K3 bf16 decode GEMM rows without hipblaslt
aryatsc Oct 29, 2026
2147087
Drop the N=6144, K=7168 a8w8 rows and sync the Kimi-K3 untuned lists
aryatsc Oct 30, 2026
77eb4f9
Merge branch 'main' into devin/k3-gemm-tables-1793030216
aryatsc Oct 30, 2026
9c953ae
Merge branch 'main' of https://git-manager.devin.ai/proxy/github.com/…
aryatsc Nov 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
53 changes: 53 additions & 0 deletions aiter/configs/model_configs/a8w8_bpreshuffle_tuned_gemm_kimik3.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1162,6 +1162,59 @@ gfx950,256,15488,7168,1536,torch.float8_e4m3fn,flydsl,1000034,0,174.6294,flydsl_
gfx950,256,15744,7168,1536,torch.float8_e4m3fn,flydsl,473,0,172.9698,flydsl_bpreshuflle_96x256x128_F8_F8_B16_1x1x4x2_default,2004.31,1508.35,0.0
gfx950,256,16000,7168,1536,torch.float8_e4m3fn,flydsl,1000034,0,175.6231,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_4x4,2006.12,1508.7,0.0
gfx950,256,16256,7168,1536,torch.float8_e4m3fn,flydsl,1000031,0,184.4351,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_3x4,1940.84,1458.64,0.0
gfx950,256,8,768,7168,torch.float8_e4m3fn,flydsl,2314,0,6.9805,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x0x0x2_default,12.62,798.61,0.0
gfx950,256,16,768,7168,torch.float8_e4m3fn,flydsl,2390,0,7.3195,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x1x0x1_default,24.07,771.13,0.0
gfx950,256,24,768,7168,torch.float8_e4m3fn,flydsl,2406,0,6.9219,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x1_default,38.17,825.48,0.0
gfx950,256,32,768,7168,torch.float8_e4m3fn,flydsl,2406,0,7.1356,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x1_default,49.38,810.52,0.0
gfx950,256,40,768,7168,torch.float8_e4m3fn,flydsl,2394,0,7.2564,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x1x4x1_default,60.69,806.62,0.0
gfx950,256,48,768,7168,torch.float8_e4m3fn,flydsl,2374,0,7.1809,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x0x0x1_default,73.6,824.8,0.0
gfx950,256,56,768,7168,torch.float8_e4m3fn,flydsl,2374,0,7.0025,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x0x0x1_default,88.05,855.75,0.0
gfx950,256,64,768,7168,torch.float8_e4m3fn,flydsl,2406,0,7.1556,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x1_default,98.47,847.18,0.0
gfx950,256,56,3584,7168,torch.float8_e4m3fn,flydsl,2379,0,9.802,flydsl_bpreshuflle_32x32x1024_F8_F8_B16_0x0x4x1_default,293.54,2702.8,0.0
gfx950,256,56,7168,512,torch.float8_e4m3fn,flydsl,9,0,3.2712,flydsl_bpreshuflle_32x64x256_F8_F8_B16_0x0x0x2_default,125.66,1376.12,0.0
gfx950,256,40,7168,1792,torch.float8_e4m3fn,flydsl,906,0,6.7691,flydsl_bpreshuflle_16x128x256_F8_F8_B16_1x3x4x2_default,151.81,1992.9,0.0
gfx950,256,1,6016,7168,torch.float8_e4m3fn,flydsl,2374,0,8.9508,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x0x0x1_default,9.64,4819.89,0.0
gfx950,256,2,6016,7168,torch.float8_e4m3fn,flydsl,2406,0,9.0277,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x1_default,19.11,4780.96,0.0
gfx950,256,4,6016,7168,torch.float8_e4m3fn,flydsl,2406,0,9.154,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x1_default,37.69,4719.19,0.0
gfx950,256,8,6016,7168,torch.float8_e4m3fn,flydsl,2330,0,9.1235,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x1x0x2_default,75.62,4743.39,0.0
gfx950,256,16,6016,7168,torch.float8_e4m3fn,flydsl,2346,0,9.7707,flydsl_bpreshuflle_16x32x1024_F8_F8_B16_0x2x0x2_default,141.23,4444.91,0.0
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
Comment thread
demonsan marked this conversation as resolved.
Resolved
gfx950,256,56,6016,7168,torch.float8_e4m3fn,flydsl,566,0,14.3068,flydsl_bpreshuflle_32x64x512_F8_F8_B16_0x2x4x2_default,337.58,3089.29,0.0
gfx950,256,64,6016,7168,torch.float8_e4m3fn,flydsl,566,0,14.5979,flydsl_bpreshuflle_32x64x512_F8_F8_B16_0x2x4x2_default,378.12,3038.21,0.0
gfx950,256,72,6016,7168,torch.float8_e4m3fn,flydsl,2061,0,18.9864,flydsl_bpreshuflle_16x64x512_F8_F8_B16_1x3x4x1_default,327.06,2344.05,0.0
gfx950,256,80,6016,7168,torch.float8_e4m3fn,flydsl,2264,0,19.6848,flydsl_bpreshuflle_16x64x512_F8_F8_B16_1x4x4x1_default,350.51,2268.69,0.0
gfx950,256,88,6016,7168,torch.float8_e4m3fn,flydsl,693,0,19.8312,flydsl_bpreshuflle_48x64x256_F8_F8_B16_1x2x4x2_default,382.71,2259.69,0.0
gfx950,256,96,6016,7168,torch.float8_e4m3fn,flydsl,1352,0,20.1734,flydsl_bpreshuflle_32x64x512_F8_F8_B16_1x0x4x1_default,410.42,2228.97,0.0
gfx950,256,104,6016,7168,torch.float8_e4m3fn,cktile,179,0,20.4731,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x4_64x64x256_1x4x1_16x16x128_default,438.11,2203.84,0.0
gfx950,256,112,6016,7168,torch.float8_e4m3fn,cktile,179,0,20.5839,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x4_64x64x256_1x4x1_16x16x128_default,469.27,2199.44,0.0
gfx950,256,120,6016,7168,torch.float8_e4m3fn,cktile,179,0,20.6189,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x4_64x64x256_1x4x1_16x16x128_default,501.94,2203.16,0.0
gfx950,256,128,6016,7168,torch.float8_e4m3fn,ck,114,0,20.7988,a8w8_bpreshuffle_256x64x64x256_16x16_16x16_16x16x1_16x16x1_1x32x1x8_8x8x1_2x1_intrawave_v3,530.77,2191.49,0.0
gfx950,256,192,6016,7168,torch.float8_e4m3fn,flydsl,204,0,23.1267,flydsl_bpreshuflle_48x128x256_F8_F8_B16_1x0x4x2_default,716.02,2024.03,0.0
gfx950,256,256,6016,7168,torch.float8_e4m3fn,flydsl,210,0,25.3278,flydsl_bpreshuflle_64x128x256_F8_F8_B16_1x0x4x2_default,871.72,1896.65,0.0
gfx950,256,384,6016,7168,torch.float8_e4m3fn,flydsl,470,0,30.0557,flydsl_bpreshuflle_96x128x256_F8_F8_B16_1x1x4x2_default,1101.89,1680.06,0.0
gfx950,256,512,6016,7168,torch.float8_e4m3fn,flydsl,702,0,33.6756,flydsl_bpreshuflle_64x192x256_F8_F8_B16_1x2x4x2_default,1311.26,1572.45,0.0
gfx950,256,768,6016,7168,torch.float8_e4m3fn,flydsl,714,0,41.7228,flydsl_bpreshuflle_96x192x256_F8_F8_B16_1x2x4x2_default,1587.54,1386.97,0.0
gfx950,256,1024,6016,7168,torch.float8_e4m3fn,flydsl,725,0,50.8153,flydsl_bpreshuflle_128x192x128_F8_F8_B16_1x2x4x2_default,1737.97,1235.52,0.0
gfx950,256,1536,6016,7168,torch.float8_e4m3fn,flydsl,1380,0,67.3911,flydsl_bpreshuflle_96x192x256_F8_F8_B16_1x0x4x1_default,1965.73,1077.5,0.0
gfx950,256,2048,6016,7168,torch.float8_e4m3fn,flydsl,1000028,0,79.6261,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_2x4,2218.25,1035.39,0.0
gfx950,256,3072,6016,7168,torch.float8_e4m3fn,flydsl,235,0,122.9702,flydsl_bpreshuflle_128x192x128_F8_F8_B16_1x0x4x2_default,2154.55,830.32,0.0
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
Comment thread
demonsan marked this conversation as resolved.
Resolved
gfx950,256,88,6144,7168,torch.float8_e4m3fn,flydsl,914,0,16.7963,flydsl_bpreshuflle_32x64x512_F8_F8_B16_1x3x4x2_default,461.48,2723.95,0.0
gfx950,256,104,6144,7168,torch.float8_e4m3fn,cktile,179,0,20.5232,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x4_64x64x256_1x4x1_16x16x128_default,446.34,2244.47,0.0
gfx950,256,112,6144,7168,torch.float8_e4m3fn,cktile,179,0,20.5456,a8w8_bpreshuffle_cktile_0x0x8x4x1x0x0x0x0x4_64x64x256_1x4x1_16x16x128_default,480.15,2249.59,0.0
gfx950,256,120,6144,7168,torch.float8_e4m3fn,flydsl,2379,0,20.5355,flydsl_bpreshuflle_32x32x1024_F8_F8_B16_0x0x4x1_default,514.7,2258.28,0.0
gfx950,256,192,6144,7168,torch.float8_e4m3fn,flydsl,452,0,23.2252,flydsl_bpreshuflle_48x128x256_F8_F8_B16_1x1x4x2_default,728.15,2057.06,0.0
gfx950,256,384,6144,7168,torch.float8_e4m3fn,flydsl,470,0,29.9909,flydsl_bpreshuflle_96x128x256_F8_F8_B16_1x1x4x2_default,1127.77,1717.56,0.0
gfx950,256,768,6144,7168,torch.float8_e4m3fn,flydsl,224,0,43.1272,flydsl_bpreshuflle_96x192x256_F8_F8_B16_1x0x4x2_default,1568.52,1367.64,0.0
gfx950,256,1536,6144,7168,torch.float8_e4m3fn,ck,49,0,65.6459,a8w8_bpreshuffle_256x192x192x128_16x16_16x16_8x32x1_8x32x1_1x32x1x8_8x8x1_2x1_intrawave_v3,2060.93,1126.11,0.0
gfx950,256,3072,6144,7168,torch.float8_e4m3fn,ck,49,0,123.0306,a8w8_bpreshuffle_256x192x192x128_16x16_16x16_8x32x1_8x32x1_1x32x1x8_8x8x1_2x1_intrawave_v3,2199.31,843.77,0.0
gfx950,256,16384,6144,7168,torch.float8_e4m3fn,flydsl,1000028,0,485.9848,flydsl_bpreshuffle_8w_256x256x128_F8_F8_B16_2x4,2969.45,746.54,0.0
gfx950,256,16,7168,3584,torch.float8_e4m3fn,ck,10,0,7.7344,a8w8_bpreshuffle_128x16x32x512_16x16_16x16_32x4x1_32x4x1_1x16x1x8_4x4x1_1x1_intrawave_v1,106.29,3358.61,0.0
gfx950,256,72,6336,7168,torch.float8_e4m3fn,flydsl,2061,0,19.224,flydsl_bpreshuflle_16x64x512_F8_F8_B16_1x3x4x1_default,340.2,2436.79,0.0
gfx950,256,72,7168,3584,torch.float8_e4m3fn,flydsl,1092,2,10.5802,flydsl_bpreshuflle_96x64x256_F8_F8_B16_1x4x0x2_default_ks2,349.65,2550.08,0.0
Expand Down
42 changes: 42 additions & 0 deletions aiter/configs/model_configs/a8w8_bpreshuffle_untuned_gemm_kimik3.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -426,6 +426,48 @@ M,N,K,q_dtype_w
8192,3072,128,torch.float8_e4m3fn
16384,3072,128,torch.float8_e4m3fn
32768,3072,128,torch.float8_e4m3fn
1,6016,7168,torch.float8_e4m3fn
2,6016,7168,torch.float8_e4m3fn
4,6016,7168,torch.float8_e4m3fn
8,6016,7168,torch.float8_e4m3fn
16,6016,7168,torch.float8_e4m3fn
24,6016,7168,torch.float8_e4m3fn
32,6016,7168,torch.float8_e4m3fn
40,6016,7168,torch.float8_e4m3fn
48,6016,7168,torch.float8_e4m3fn
56,6016,7168,torch.float8_e4m3fn
64,6016,7168,torch.float8_e4m3fn
72,6016,7168,torch.float8_e4m3fn
80,6016,7168,torch.float8_e4m3fn
88,6016,7168,torch.float8_e4m3fn
96,6016,7168,torch.float8_e4m3fn
104,6016,7168,torch.float8_e4m3fn
112,6016,7168,torch.float8_e4m3fn
120,6016,7168,torch.float8_e4m3fn
128,6016,7168,torch.float8_e4m3fn
192,6016,7168,torch.float8_e4m3fn
256,6016,7168,torch.float8_e4m3fn
384,6016,7168,torch.float8_e4m3fn
512,6016,7168,torch.float8_e4m3fn
768,6016,7168,torch.float8_e4m3fn
1024,6016,7168,torch.float8_e4m3fn
1536,6016,7168,torch.float8_e4m3fn
2048,6016,7168,torch.float8_e4m3fn
3072,6016,7168,torch.float8_e4m3fn
4096,6016,7168,torch.float8_e4m3fn
8192,6016,7168,torch.float8_e4m3fn
16384,6016,7168,torch.float8_e4m3fn
40,6144,7168,torch.float8_e4m3fn
88,6144,7168,torch.float8_e4m3fn
104,6144,7168,torch.float8_e4m3fn
112,6144,7168,torch.float8_e4m3fn
120,6144,7168,torch.float8_e4m3fn
192,6144,7168,torch.float8_e4m3fn
384,6144,7168,torch.float8_e4m3fn
768,6144,7168,torch.float8_e4m3fn
1536,6144,7168,torch.float8_e4m3fn
3072,6144,7168,torch.float8_e4m3fn
16384,6144,7168,torch.float8_e4m3fn
16,7168,3584,torch.float8_e4m3fn
32,6336,7168,torch.float8_e4m3fn
72,6336,7168,torch.float8_e4m3fn
Expand Down
Loading
Loading