From a824b93c7b3dfa757db6ed93356003e9ad2e3de5 Mon Sep 17 00:00:00 2001 From: "jinzhen.ljz" Date: Mon, 12 Oct 2026 23:08:25 +0800 Subject: [PATCH 1/4] fix(sm90): correct WGMMA register checks and grouped tile limits Co-authored-by: Codex (GPT-6) Signed-off-by: jinzhen.ljz --- humming/config/mma.py | 42 +++++++++++------------------------ humming/tune/sm90_policies.py | 6 ++++- 2 files changed, 18 insertions(+), 30 deletions(-) diff --git a/humming/config/mma.py b/humming/config/mma.py index adbe8cc3..116d25d5 100644 --- a/humming/config/mma.py +++ b/humming/config/mma.py @@ -191,13 +191,6 @@ def get_register_budget_error(layer_config, tuning_config, use_f16_accum=False, f"{instruction_registers} registers per thread; launch budget {launch_budget:g}" ) - accumulator_registers = warp_m * warp_n / (64 if use_f16_accum else 32) - has_group_accumulator = False - if mma_type != MmaType.MXMMA and layer_config.a_dtype.num_bits < 16: - has_group_scale = layer_config.input_scale_group_size > 0 - if not has_group_scale and not layer_config.use_fused_e8m0_scale: - has_group_scale = layer_config.is_group_weight_scale or layer_config.is_block_weight_scale - has_group_accumulator = has_group_scale math_budget = min(255, launch_budget) if tuning_config.use_warp_spec: try: @@ -211,33 +204,24 @@ def get_register_budget_error(layer_config, tuning_config, use_f16_accum=False, f"launch budget {launch_budget:g}, math-thread budget {math_budget:g}" ) - # Each ordinary MMA buffer spans K=256/activation_bits, hence M/4 and N/4 - # 32-bit registers per thread. Count dequantized operands, not both copies. - buffer_registers = (warp_m + warp_n) / 4 if mma_type == MmaType.WGMMA: - # A is always in smem; raw-weight SS also keeps B in smem. - buffer_registers = 0 if layer_config.use_raw_weight else warp_n / 4 - if not layer_config.use_raw_weight and layer_config.use_packed_k_layout: - # The packed B buffer holds one N=64 WGMMA fragment across all K slabs. - buffer_registers = warp_k * layer_config.a_dtype.num_bits / 64 + return None + + accumulator_registers = warp_m * warp_n / (64 if use_f16_accum else 32) + has_group_accumulator = False + if mma_type != MmaType.MXMMA and layer_config.a_dtype.num_bits < 16: + has_group_scale = layer_config.input_scale_group_size > 0 + if not has_group_scale and not layer_config.use_fused_e8m0_scale: + has_group_scale = layer_config.is_group_weight_scale or layer_config.is_block_weight_scale + has_group_accumulator = has_group_scale + buffer_registers = (warp_m + warp_n) / 4 if has_group_accumulator: - partial_accumulator_registers = accumulator_registers * 0.5 - if mma_type == MmaType.WGMMA: - partial_accumulator_registers = max(partial_accumulator_registers, instruction_accumulators) - if layer_config.a_dtype.is_integer_type or tuning_config.wgmma_split_issue_wait: - partial_accumulator_registers = accumulator_registers - accumulator_registers += partial_accumulator_registers - if mma_type == MmaType.WGMMA and tuning_config.wgmma_split_issue_wait: - accumulator_registers = max(accumulator_registers, 2 * instruction_accumulators) - input_scale_registers = 0 - if mma_type == MmaType.WGMMA and layer_config.input_scale_group_size > 0: - input_scale_registers = math.ceil(warp_m / 16) * 4 - demand = accumulator_registers + buffer_registers + input_scale_registers + accumulator_registers *= 1.5 + demand = accumulator_registers + buffer_registers if demand >= math_budget - 8: return ( f"register budget exceeded: accumulator {accumulator_registers:g} + " - f"single-buffer {buffer_registers:g} + input scales {input_scale_registers:g} " - f"= {demand:g} must be < " + f"single-buffer {buffer_registers:g} = {demand:g} must be < " f"math-thread budget {math_budget:g} - 8 ({math_budget - 8:g})" ) return None diff --git a/humming/tune/sm90_policies.py b/humming/tune/sm90_policies.py index 873b2c99..ff0efb21 100644 --- a/humming/tune/sm90_policies.py +++ b/humming/tune/sm90_policies.py @@ -42,6 +42,10 @@ def fit_sm90_tile( has_weight_scale_accumulator = has_group_weight_scale and not layer_config.use_fused_e8m0_scale has_scaled_accumulator = layer_config.is_group_input_scale or has_weight_scale_accumulator has_group_accumulator = layer_config.a_dtype.num_bits < 16 and has_scaled_accumulator + has_packed_input_scales = layer_config.use_packed_k_layout and layer_config.is_group_input_scale + has_full_group_accumulator = layer_config.a_dtype.is_integer_type + has_full_group_accumulator |= config.get("wgmma_split_issue_wait", False) + max_group_warp_m = 80 if has_packed_input_scales and has_full_group_accumulator else 96 has_unpacked_rs = not layer_config.use_raw_weight and not layer_config.use_packed_k_layout should_widen_rs = has_unpacked_rs and has_group_accumulator and not use_warp_spec and warp_n < 32 if mma_type == MmaType.WGMMA and should_widen_rs and layer_config.shape_n % 128 == 0: @@ -65,7 +69,7 @@ def fit_sm90_tile( if layer_config.a_dtype.num_bits == 16 and use_warp_spec and warp_n <= 32: max_warp_m = 192 if has_group_accumulator: - max_warp_m = 160 if warp_n <= 16 else 96 if warp_n <= 32 else 48 + max_warp_m = 160 if warp_n <= 16 else max_group_warp_m if warp_n <= 32 else 48 if has_unpacked_rs: max_warp_m = 128 if warp_n <= 16 else 80 if warp_n <= 32 else 48 if layer_config.a_dtype.is_integer_type: From 592b924fa94b6147541bdd5f3e496e10f9175a55 Mon Sep 17 00:00:00 2001 From: "jinzhen.ljz" Date: Tue, 13 Oct 2026 01:56:17 +0800 Subject: [PATCH 2/4] fix(sm90): reserve sufficient registers for WGMMA RS issue Co-authored-by: Codex (GPT-6) Signed-off-by: jinzhen.ljz --- humming/config/mma.py | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/humming/config/mma.py b/humming/config/mma.py index 116d25d5..77dac209 100644 --- a/humming/config/mma.py +++ b/humming/config/mma.py @@ -174,11 +174,10 @@ def get_register_budget_error(layer_config, tuning_config, use_f16_accum=False, return None # UMMA accumulators use TMEM and have separate resource checks. if mma_type == MmaType.WGMMA: - # One instruction spans project N=64 across four warps. ptxas needs - # 26 extra registers for SS, or 30 for RS, beyond its accumulator tuple. - # setmaxnreg can raise the runtime math budget, but not this compile limit. + # ptxas needs 26 (SS) or 32 (RS) registers beyond one N=64 accumulator tuple. + # setmaxnreg cannot relax this per-instruction compile limit. instruction_accumulators = warp_m // (4 if use_f16_accum else 2) - instruction_overhead = 26 if layer_config.use_raw_weight else 30 + instruction_overhead = 26 if layer_config.use_raw_weight else 32 instruction_registers = instruction_accumulators + instruction_overhead if layer_config.use_packed_k_layout and tuning_config.wgmma_split_issue_wait: # Split packed-K keeps all RS K slabs live across prefetch. The From 7f22a9003efac561ccfde7d53dd3ceec803f80a5 Mon Sep 17 00:00:00 2001 From: "jinzhen.ljz" Date: Tue, 13 Oct 2026 01:56:45 +0800 Subject: [PATCH 3/4] ci(h200): use nvcc for sampled kernel tests Co-authored-by: Codex (GPT-6) Signed-off-by: jinzhen.ljz --- .buildkite/pipeline.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index b2b54ebc..41c26234 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -160,4 +160,4 @@ steps: queue: h200_35gb concurrency_group: humming/h200-sampled depends_on: approve-h200-sampled-tests - command: HUMMING_TEST_TUNING_SEED=20261005 HUMMING_TEST_TUNING_SOURCE=sampled32 bash .buildkite/run-tests.sh tests/kernels/humming/ + command: HUMMING_COMPILER=nvcc HUMMING_TEST_TUNING_SEED=20261005 HUMMING_TEST_TUNING_SOURCE=sampled32 bash .buildkite/run-tests.sh tests/kernels/humming/ From 025fc4d315e709f40c9907670f4fa3d63f05e2fd Mon Sep 17 00:00:00 2001 From: "jinzhen.ljz" Date: Tue, 13 Oct 2026 10:21:07 +0800 Subject: [PATCH 4/4] revert(ci): restore default compiler for H200 sampled tests This reverts commit 7f22a90. Signed-off-by: jinzhen.ljz --- .buildkite/pipeline.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index 41c26234..b2b54ebc 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -160,4 +160,4 @@ steps: queue: h200_35gb concurrency_group: humming/h200-sampled depends_on: approve-h200-sampled-tests - command: HUMMING_COMPILER=nvcc HUMMING_TEST_TUNING_SEED=20261005 HUMMING_TEST_TUNING_SOURCE=sampled32 bash .buildkite/run-tests.sh tests/kernels/humming/ + command: HUMMING_TEST_TUNING_SEED=20261005 HUMMING_TEST_TUNING_SOURCE=sampled32 bash .buildkite/run-tests.sh tests/kernels/humming/