Skip to content

fix(sm90): correct WGMMA register checks and grouped tile limits #138

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

Merged
jinzhen-lin merged 4 commits into from
Oct 13, 2026
Merged

fix(sm90): correct WGMMA register checks and grouped tile limits #138

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
a824b93
fix(sm90): correct WGMMA register checks and grouped tile limits
jinzhen-lin Oct 12, 2026
592b924
fix(sm90): reserve sufficient registers for WGMMA RS issue
jinzhen-lin Oct 12, 2026
7f22a90
ci(h200): use nvcc for sampled kernel tests
jinzhen-lin Oct 12, 2026
025fc4d
revert(ci): restore default compiler for H200 sampled tests
jinzhen-lin Oct 13, 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
49 changes: 16 additions & 33 deletions humming/config/mma.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -191,13 +190,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:
Expand All @@ -211,33 +203,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
Expand Down
6 changes: 5 additions & 1 deletion humming/tune/sm90_policies.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand All @@ -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:
Expand Down
Loading