Skip to content

Commit a9c758d

Browse files
apstenku123apstenku123
committed
refactor: retire legacy Path B MSL bypass and consolidate top-k selection into TileLang Path C owner-output routes
1 parent 86c487e commit a9c758d

63 files changed

Lines changed: 7475 additions & 6172 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎bench/tilelang_ports/fp8_msl_kernels.json‎

Lines changed: 30 additions & 183 deletions
Original file line numberDiff line numberDiff line change
@@ -1,173 +1,20 @@
11
{
22
"schema_version": 1,
33
"scope": "local_only",
4-
"kernel": "fp8_msl_kernels",
5-
"license_notice": "Vendored from AppMana/mps-fp8-for-torch-and-comfyui-python-package (commit a902571e, Apache 2.0) and audiohacking/fp8-mps-metal (commit d4fbd40c, MIT).",
4+
"kernel": "fp8_reference_helpers",
5+
"license_notice": "Historical direct-MSL sources were vendored from AppMana/mps-fp8-for-torch-and-comfyui-python-package (commit a902571e, Apache 2.0) and audiohacking/fp8-mps-metal (commit d4fbd40c, MIT); current helpers are pure MLX references.",
66
"hardware_label": "Davids-Mac-Studio.local",
77
"platform": {
88
"system": "Darwin",
99
"machine": "arm64",
1010
"python_version": "3.13.12",
11-
"mlx_version": "0.31.1"
11+
"mlx_version": "0.32.0.dev20260514+d168ca5ca"
1212
},
1313
"metal_status": {
14-
"available": true,
15-
"reason": "Vendored FP8 e4m3fn MSL kernels (LUT-based decode + integer-bit encode) are compiled and ready to dispatch."
14+
"available": false,
15+
"reason": "FP8 direct-MSL Path B is retired for production cleanup: this module no longer constructs or dispatches vendored Metal kernels. Use fp8_matmul_path_c.py or fp8_vecmat_path_c.py for the TileLang/tvm-ffi owner-output route, or these helpers only as pure-MLX reference oracles."
1616
},
1717
"shapes": [
18-
{
19-
"shape_label": "square_4096",
20-
"M": 4096,
21-
"N": 4096,
22-
"K": 4096,
23-
"fp8_input_bytes": 33554432,
24-
"fp16_input_bytes": 67108864,
25-
"memory_savings_pct": 50.0,
26-
"rows": [
27-
{
28-
"label": "square_4096/half_to_fp8 (A 4096x4096)",
29-
"median_ms": 1.1150420177727938,
30-
"min_ms": 1.015292014926672,
31-
"max_ms": 2.0716250874102116,
32-
"iters": 10,
33-
"warmup": 3
34-
},
35-
{
36-
"label": "square_4096/fp8_to_half (A 4096x4096)",
37-
"median_ms": 0.6388330366462469,
38-
"min_ms": 0.5903330165892839,
39-
"max_ms": 1.3812498655170202,
40-
"iters": 10,
41-
"warmup": 3
42-
},
43-
{
44-
"label": "square_4096/fp8_scaled_matmul (4096x4096 @ 4096x4096)",
45-
"median_ms": 228.22708287276328,
46-
"min_ms": 225.1199579332024,
47-
"max_ms": 232.51062515191734,
48-
"iters": 10,
49-
"warmup": 3
50-
},
51-
{
52-
"label": "square_4096/mx_matmul_fp16 (4096x4096 @ 4096x4096)",
53-
"median_ms": 9.78454202413559,
54-
"min_ms": 9.542542044073343,
55-
"max_ms": 12.967708986252546,
56-
"iters": 10,
57-
"warmup": 3
58-
},
59-
{
60-
"label": "square_4096/fp8_scaled_vecmat (1x4096 @ 4096x4096)",
61-
"median_ms": 0.2628329675644636,
62-
"min_ms": 0.181916868314147,
63-
"max_ms": 0.43816608376801014,
64-
"iters": 10,
65-
"warmup": 3
66-
}
67-
]
68-
},
69-
{
70-
"shape_label": "tall_skinny",
71-
"M": 2048,
72-
"N": 512,
73-
"K": 4096,
74-
"fp8_input_bytes": 10485760,
75-
"fp16_input_bytes": 20971520,
76-
"memory_savings_pct": 50.0,
77-
"rows": [
78-
{
79-
"label": "tall_skinny/half_to_fp8 (A 2048x4096)",
80-
"median_ms": 0.6746661383658648,
81-
"min_ms": 0.5720830522477627,
82-
"max_ms": 0.9097501169890165,
83-
"iters": 10,
84-
"warmup": 3
85-
},
86-
{
87-
"label": "tall_skinny/fp8_to_half (A 2048x4096)",
88-
"median_ms": 0.42091705836355686,
89-
"min_ms": 0.3674591425806284,
90-
"max_ms": 0.5260000471025705,
91-
"iters": 10,
92-
"warmup": 3
93-
},
94-
{
95-
"label": "tall_skinny/fp8_scaled_matmul (2048x4096 @ 4096x512)",
96-
"median_ms": 17.41487509571016,
97-
"min_ms": 17.076958902180195,
98-
"max_ms": 17.594290897250175,
99-
"iters": 10,
100-
"warmup": 3
101-
},
102-
{
103-
"label": "tall_skinny/mx_matmul_fp16 (2048x4096 @ 4096x512)",
104-
"median_ms": 0.8477920200675726,
105-
"min_ms": 0.810291850939393,
106-
"max_ms": 0.9760840330272913,
107-
"iters": 10,
108-
"warmup": 3
109-
},
110-
{
111-
"label": "tall_skinny/fp8_scaled_vecmat (1x4096 @ 4096x512)",
112-
"median_ms": 0.16604107804596424,
113-
"min_ms": 0.1437910832464695,
114-
"max_ms": 0.21087494678795338,
115-
"iters": 10,
116-
"warmup": 3
117-
}
118-
]
119-
},
120-
{
121-
"shape_label": "fat_K",
122-
"M": 512,
123-
"N": 512,
124-
"K": 8192,
125-
"fp8_input_bytes": 8388608,
126-
"fp16_input_bytes": 16777216,
127-
"memory_savings_pct": 50.0,
128-
"rows": [
129-
{
130-
"label": "fat_K/half_to_fp8 (A 512x8192)",
131-
"median_ms": 0.3708750009536743,
132-
"min_ms": 0.34129200503230095,
133-
"max_ms": 0.42716716416180134,
134-
"iters": 10,
135-
"warmup": 3
136-
},
137-
{
138-
"label": "fat_K/fp8_to_half (A 512x8192)",
139-
"median_ms": 0.30683400109410286,
140-
"min_ms": 0.2810410223901272,
141-
"max_ms": 0.6989999674260616,
142-
"iters": 10,
143-
"warmup": 3
144-
},
145-
{
146-
"label": "fat_K/fp8_scaled_matmul (512x8192 @ 8192x512)",
147-
"median_ms": 8.045291993767023,
148-
"min_ms": 7.845416897907853,
149-
"max_ms": 8.394957985728979,
150-
"iters": 10,
151-
"warmup": 3
152-
},
153-
{
154-
"label": "fat_K/mx_matmul_fp16 (512x8192 @ 8192x512)",
155-
"median_ms": 0.5070420447736979,
156-
"min_ms": 0.4676659591495991,
157-
"max_ms": 0.6057089194655418,
158-
"iters": 10,
159-
"warmup": 3
160-
},
161-
{
162-
"label": "fat_K/fp8_scaled_vecmat (1x8192 @ 8192x512)",
163-
"median_ms": 0.17433403991162777,
164-
"min_ms": 0.15616579912602901,
165-
"max_ms": 0.2648751251399517,
166-
"iters": 10,
167-
"warmup": 3
168-
}
169-
]
170-
},
17118
{
17219
"shape_label": "tiny_smoke",
17320
"M": 64,
@@ -179,43 +26,43 @@
17926
"rows": [
18027
{
18128
"label": "tiny_smoke/half_to_fp8 (A 64x128)",
182-
"median_ms": 0.15249988064169884,
183-
"min_ms": 0.12258300557732582,
184-
"max_ms": 0.3139160107821226,
185-
"iters": 10,
186-
"warmup": 3
29+
"median_ms": 0.1433750003343448,
30+
"min_ms": 0.14179099525790662,
31+
"max_ms": 0.16500000492669642,
32+
"iters": 5,
33+
"warmup": 2
18734
},
18835
{
18936
"label": "tiny_smoke/fp8_to_half (A 64x128)",
190-
"median_ms": 0.15237508341670036,
191-
"min_ms": 0.12087495997548103,
192-
"max_ms": 0.19095814786851406,
193-
"iters": 10,
194-
"warmup": 3
37+
"median_ms": 0.13241700071375817,
38+
"min_ms": 0.131665998196695,
39+
"max_ms": 0.13666599988937378,
40+
"iters": 5,
41+
"warmup": 2
19542
},
19643
{
19744
"label": "tiny_smoke/fp8_scaled_matmul (64x128 @ 128x64)",
198-
"median_ms": 0.1487499102950096,
199-
"min_ms": 0.13329205103218555,
200-
"max_ms": 0.2335000317543745,
201-
"iters": 10,
202-
"warmup": 3
45+
"median_ms": 0.1732500022626482,
46+
"min_ms": 0.16579199291300029,
47+
"max_ms": 0.17716600268613547,
48+
"iters": 5,
49+
"warmup": 2
20350
},
20451
{
20552
"label": "tiny_smoke/mx_matmul_fp16 (64x128 @ 128x64)",
206-
"median_ms": 0.17416616901755333,
207-
"min_ms": 0.13875006698071957,
208-
"max_ms": 0.2495830412954092,
209-
"iters": 10,
210-
"warmup": 3
53+
"median_ms": 0.15145799989113584,
54+
"min_ms": 0.14766700041946024,
55+
"max_ms": 0.1525410043541342,
56+
"iters": 5,
57+
"warmup": 2
21158
},
21259
{
21360
"label": "tiny_smoke/fp8_scaled_vecmat (1x128 @ 128x64)",
214-
"median_ms": 0.1582079567015171,
215-
"min_ms": 0.12612505815923214,
216-
"max_ms": 0.24137482978403568,
217-
"iters": 10,
218-
"warmup": 3
61+
"median_ms": 0.1325839984929189,
62+
"min_ms": 0.13112500164425,
63+
"max_ms": 0.14645800547441468,
64+
"iters": 5,
65+
"warmup": 2
21966
}
22067
]
22168
}

0 commit comments

Comments
 (0)