[CI] Split (B200) Miscellaneous Kernels into mHC, FLA Ops and Misc named jobs (#58609)
Signed-off-by: Thang Nguyen <thangnguyenvn647@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com>
Thang Nguyen committed
Sep 26, 2026 at 07:10 UTC
2f41e002351188284885282a374a16e6c0136d4d
11 files changed
+31
-6
.buildkite/test_areas/kernels.yaml
+29
-4
@@ -144,6 +144,32 @@ steps:
144
- pytest -v -s kernels/test_fused_deepseek_v4_*.py
145
- pytest -v -s models/test_deepseek_v4_mega_moe.py
146
147
+# Own job so mHC never shares a process with FlashInfer's import-time
148
+# ptxas override.
149
+- label: ":nvidia: (B200) mHC Kernels"
150
+ key: kernels-mhc-test-b200
151
+ timeout_in_minutes: 25
152
+ device: b200-k8s
153
+ source_file_dependencies:
154
+ - csrc/
155
+ - vllm/
156
+ - "!vllm/distributed/kv_transfer/"
157
+ - tests/kernels/
158
+ commands:
159
+ - pytest -v -s kernels/mhc
160
+
161
+- label: ":nvidia: (B200) FLA Ops Kernels"
162
+ key: kernels-fla-ops-test-b200
163
+ timeout_in_minutes: 25
164
+ device: b200-k8s
165
+ source_file_dependencies:
166
+ - csrc/
167
+ - vllm/
168
+ - "!vllm/distributed/kv_transfer/"
169
+ - tests/kernels/
170
+ commands:
171
+ - pytest -v -s kernels/fla
172
+
173
# Catch-all for test files at the tests/kernels root. This job collects
174
# the whole root so new files are wired by default.
175
# Files with dedicated jobs elsewhere in this file are excluded via --ignore
@@ -151,7 +177,7 @@ steps:
177
# their own jobs / Kernels (B200)).
178
- label: ":nvidia: (B200) Miscellaneous Kernels"
179
key: kernels-root-misc-test-b200
154
- timeout_in_minutes: 45
180
+ timeout_in_minutes: 30
181
device: b200-k8s
182
source_file_dependencies:
183
- csrc/
@@ -159,14 +185,14 @@ steps:
185
- "!vllm/distributed/kv_transfer/"
186
- tests/kernels/
187
commands:
162
- # Isolate mHC from FlashInfer's import-time ptxas override.
163
- - pytest -v -s kernels/test_mhc_kernels.py
188
- pytest -v -s kernels/
189
--ignore=kernels/attention
190
--ignore=kernels/core
191
+ --ignore=kernels/fla
192
--ignore=kernels/helion
193
--ignore=kernels/ir
194
--ignore=kernels/mamba
195
+ --ignore=kernels/mhc
196
--ignore=kernels/moe
197
--ignore=kernels/quantization
198
--ignore=kernels/test_concat_mla_q.py
@@ -176,7 +202,6 @@ steps:
202
--ignore=kernels/test_kda.py
203
--ignore=kernels/test_bf16x3_router_gemm_cutedsl.py
204
--ignore=kernels/test_ll_bf16_gemm.py
179
- --ignore=kernels/test_mhc_kernels.py
205
--ignore=kernels/test_shuffle_rows.py
206
# BROKEN on main, pending kernel fixes (B200):
207
# test_shuffle_rows.py (1: test_shuffle_rows_edge_cases)
tests/kernels/fla/__init__.py
tests/kernels/fla/test_fla_layernorm_guard.py
renamed
tests/kernels/fla/test_fused_gdn_post_conv.py
renamed
tests/kernels/fla/test_fused_recurrent_packed_decode.py
renamed
tests/kernels/fla/test_fused_sigmoid_gating_delta_rule.py
renamed
tests/kernels/mhc/__init__.py
tests/kernels/mhc/test_mhc_jit_warmup.py
renamed
tests/kernels/mhc/test_mhc_kernels.py
renamed
tests/kernels/mhc/test_mhc_tilelang_jit.py
renamed
+1
-1
@@ -108,7 +108,7 @@ def test_deepseek_v4_import_and_jit_monitor_do_not_hijack_hip_symbols() -> None:
108
# Both claims are about process-global state, `sys.modules` and the symbol
109
# table, and a sibling test legitimately imports TileLang to exercise those
110
# kernels, so the checks only mean something in an interpreter of their own.
111
- script = Path(__file__).parent / "scripts" / "check_no_tilelang_hijack.py"
111
+ script = Path(__file__).parents[1] / "scripts" / "check_no_tilelang_hijack.py"
112
result = subprocess.run(
113
[sys.executable, str(script)],
114
capture_output=True,
vllm/model_executor/kernels/mhc/cpu.py
+1
-1
@@ -81,7 +81,7 @@ def hc_head_fused_cpu(
81
hc_eps: float,
82
) -> torch.Tensor:
83
"""CPU-ported HC head reduction (see `test_hc_head_cpu` in
84
- tests/kernels/test_mhc_kernels.py for the eager reference this is tested
84
+ tests/kernels/mhc/test_mhc_kernels.py for the eager reference this is tested
85
against).
86
87
The ported kernel's C++ signature takes `(hc_eps, norm_eps)` -- the