heuristic: cap CTAs at 2 for large-M grouped-input-scale GEMMs on H20

[guzekai01]

Summary

On H20, the grouped-input-scale mainloop requires additional live accumulator and scale state. With num_ctas_per_sm=3, __launch_bounds__ limits this 128-thread kernel to about 170 registers per thread, causing the compiler to spill to local memory.

Cap num_ctas_per_sm at 2 for 8-bit grouped-input-scale, non-grouped-masked GEMMs with shape_m >= 6144. The cap is applied after the existing shape and pipeline decisions, so all other config fields remain unchanged.

For the DSV4-Pro w13 shape (m=12288, n=6144, k=7168, E=48):

metric	CTA3 (before)	CTA2 (after)
benchmark latency	5.612 ms	4.247 ms (-24.3%)
registers/thread	168	236
local loads/stores	34.2M / 19.2M	0 / 0
spilling requests	53.6M	0

The additional occupancy from CTA3 does not offset the spill cost at large M. Small-M and grouped-masked cases are left unchanged because they still benefit from the extra CTA.

[gemini-code-assist]

Code Review

This pull request introduces a cap on num_ctas_per_sm to at most 2 under specific conditions in humming/tune/sm90_h20.py. The reviewer pointed out two critical issues: first, capping num_ctas_per_sm leaves num_warps with an outdated value, which can hinder downstream optimizations; second, meta.input_scale_group_size can be None, potentially causing a TypeError during comparison. A code suggestion was provided to resolve both issues.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

[gemini-code-assist]

There are two issues with this block:

1. Outdated num_warps: num_warps is calculated on line 155 using the uncapped num_ctas_per_sm. Capping num_ctas_per_sm here leaves num_warps with an outdated, higher value. This can prevent downstream optimizations (such as enabling TMA, warp specialization, and mbarriers on line 236) from being applied because they check num_warps <= 8. Recalculating num_warps when num_ctas_per_sm is capped resolves this.
2. Defensive check for None: meta.input_scale_group_size can be None (as seen in other parts of the codebase where it is checked with an or fallback). Comparing it directly with > 0 can raise a TypeError. Using (meta.input_scale_group_size or 0) > 0 is safer and more robust.

        if (
            meta.a_dtype.num_bits == 8
            and (meta.input_scale_group_size or 0) > 0
            and gemm_type != GemmType.GROUPED_MASKED
            and shape_m >= 6144
        ):
            if num_ctas_per_sm > 2:
                num_ctas_per_sm = 2
                num_warps = num_warps_m * num_warps_n * num_warps_k * num_ctas_per_sm