Skip to content

[autotuner][cute] add formula-based tcgen05 matmul seed heuristic - #3152

Draft
calebmkim wants to merge 1 commit into
calebmkim/stack/35from
calebmkim/stack/36
Draft

[autotuner][cute] add formula-based tcgen05 matmul seed heuristic#3152
calebmkim wants to merge 1 commit into
calebmkim/stack/35from
calebmkim/stack/36

Conversation

@calebmkim

@calebmkim calebmkim commented Jul 24, 2026

Copy link
Copy Markdown
Contributor

Stacked PRs:


[autotuner][cute] add formula-based tcgen05 matmul seed heuristic

Replaces the special-case cluster_m=2 seed producers with one shape-aware analytical
formula (cute analog of the Triton B200 formula heuristic #3007):
f(M,K,N,dtype,epilogue,num_sm,smem_budget) -> Config, promoted to the no-autotune
default. Covers configs the old producers structurally could not emit: cluster_m=1
decode (M<=128, previously NO seed) and deep-AB compute (ab=6 vs the old ab=3).

Depends on PRs 1-3 below it (they widen the search surface / unblock injection); this
PR is the config PRODUCER on top of that surface.

STATUS: this is a STARTING POINT, not a finished heuristic. It is a first, deliberately
simple analytical formula (~5 knobs) that already covers regimes the old producers
could not and beats them on quality (below), but it is expected to be iterated on:
the classifier thresholds (e.g. the wave-fill / _MANY_WAVE crossovers), the l2_groupings
rule, and coverage of more epilogues/shapes are all approximations with known gaps (see
the F1 limitation below). Intent is to land a maintainable baseline that replaces the
special-case producers and then improve it incrementally, not to ship a final tuning oracle.

What this PR is responsible for (measured this run, single B200, cold-L2 cudagraph):

[A] Config quality -- the config the formula EMITS vs the old producers / an external
reference, benched directly (no autotuner):

  • geomean(seed / old-producer) = 1.091 over the 23 curriculum shapes where an old
    producer fired, 0 regressions (0 shapes >2% slower). The old producers emit NO
    config on the other 12 shapes (cluster_m=1 decode etc.) -- the "A1 gap" this
    formula closes.
  • geomean(seed / external ref) = 1.035 (curriculum), 1.083 (held-out fp16, 3 shapes
    with a ref) -- held-out >= curriculum, so no overfitting.
  • accuracy 35/35 PASS (curriculum core; 48/48 incl. the FFI alt-seed arm; 4/4 fp16).
  • FFI alt-seed KEPT: 3 curriculum shapes show a >2% win of the FFI explicit_epi_tile
    topology over the DEFAULT-layout core -- residual_add 4096^3 +54.0%,
    bias_residual_gelu 8192^3 +50.6%, silu 4096^3 +9.5% (the source-C/residual family
    where the core caps ab at 2 and the FFI ring recovers it). The autotuner keeps
    whichever of {core, FFI alt} wins.

[A] No-autotune default -- what a user gets with autotuning OFF (default_config()),
the promoted formula config vs the raw fragment default:

  • 64x8192x8192 fp8 321.9 vs 29.3 TF/s (~11x); 512x4096x4096 bf16 697.2 vs 557.5;
    2048x4096x4096 bf16 1340.5 vs 1289.8. This is the primary user-visible win of the
    seed and does not depend on running the search.

[B] Effect on the autotune SEARCH (formula seed ON vs OFF, both on this full stack):

  • Convergence: on 64x8192x8192 fp8 the seed lands the search AT the answer -- the
    first benchmarked config is ~1.05x the final best (seeded) vs ~4.77x (unseeded),
    reproducible across two reps. Both arms ultimately converge to the SAME winner
    ([64,64,128] cm1 ab12), re-benched cold-L2 to an identical 321.9 TF/s (any
    end-of-search delta in the tables is the autotuner's noisy in-process timer, not a
    real difference -- the cold-L2 re-bench is the trustworthy comparison).
  • Ceiling: given PRs 1-3, the unseeded search on this stack already REACHES the good
    cluster_m=2 configs (e.g. 512x4096x4096 bf16 base and treat both reach
    [256,128,128] cm2 ab4 = 697.2). That reachability is PR3's contribution (widening
    the search space), which the seed-OFF arm confirms -- NOT a win attributable to
    this seed. The formula's marginal search-ceiling effect is small (a soft +3.9% on
    2048x4096x4096 bf16, where the seeded search reached the deep-AB [256,256,64] ab6
    tile in 1 of 2 reps and the unseeded did not). The seed's real value is the
    no-autotune default and convergence above, not moving the search ceiling.

Prerequisite edits bundled here (needed for the formula to emit valid configs):

  • cute_mma.py: admit a rank-2 exact-shape residual aux on the explicit-epi-tile path
    (broadcast_axis in (1, None)) for the bias_residual_gelu / residual_add family --
    without it the FFI alt-seed's residual config is rejected (it drives the +50% wins).
  • tcgen05_config.py: extend the 16-bit deep-AB (ab>3) SMEM-clamped admission from
    cluster_m=2 to cluster_m=1 DEFAULT layout, so the formula's bf16 decode tile
    ([64,32,256] ab4 = 196608 B) validates. max_ab_stages_that_fit still rejects the
    real cm1 256^2 overflow (clamped to ab<=3).
  • _gelu_tanh_approx.py: lower cute.math.tanh with fastmath=True (hw tanh.approx.f32),
    recovering ~14% on the fused GEMM+GELU epilogue.

Also pins the promoted DEFAULT-layout seed's layout/ffi/epi-tile keys explicitly: on a
direct-entry-eligible shape the fragment defaults are FFI-biased, so a promoted seed
that omitted them merged into an invalid explicit_epi_tile+ffi hybrid and raised
InvalidConfig in the no-autotune / baseline compile path.

Known limitation, NOT fixed (out of scope): fp8 512x2048x2048 -- the formula emits
l2_groupings=[4] where the old producer emits [1]. Measured this run it is a near-tie
(seed/old 0.997), but the l2_groupings divergence is real; left for a follow-up.

Co-Authored-By: Claude Opus 4.8 (1M context) noreply@anthropic.com

Replaces the special-case cluster_m=2 seed producers with one shape-aware analytical
formula (cute analog of the Triton B200 formula heuristic #3007):
f(M,K,N,dtype,epilogue,num_sm,smem_budget) -> Config, promoted to the no-autotune
default. Covers configs the old producers structurally could not emit: cluster_m=1
decode (M<=128, previously NO seed) and deep-AB compute (ab=6 vs the old ab=3).

Depends on PRs 1-3 below it (they widen the search surface / unblock injection); this
PR is the config PRODUCER on top of that surface.

STATUS: this is a STARTING POINT, not a finished heuristic. It is a first, deliberately
simple analytical formula (~5 knobs) that already covers regimes the old producers
could not and beats them on quality (below), but it is expected to be iterated on:
the classifier thresholds (e.g. the wave-fill / _MANY_WAVE crossovers), the l2_groupings
rule, and coverage of more epilogues/shapes are all approximations with known gaps (see
the F1 limitation below). Intent is to land a maintainable baseline that replaces the
special-case producers and then improve it incrementally, not to ship a final tuning oracle.

What this PR is responsible for (measured this run, single B200, cold-L2 cudagraph):

[A] Config quality -- the config the formula EMITS vs the old producers / an external
    reference, benched directly (no autotuner):
  * geomean(seed / old-producer) = 1.091 over the 23 curriculum shapes where an old
    producer fired, 0 regressions (0 shapes >2% slower). The old producers emit NO
    config on the other 12 shapes (cluster_m=1 decode etc.) -- the "A1 gap" this
    formula closes.
  * geomean(seed / external ref) = 1.035 (curriculum), 1.083 (held-out fp16, 3 shapes
    with a ref) -- held-out >= curriculum, so no overfitting.
  * accuracy 35/35 PASS (curriculum core; 48/48 incl. the FFI alt-seed arm; 4/4 fp16).
  * FFI alt-seed KEPT: 3 curriculum shapes show a >2% win of the FFI explicit_epi_tile
    topology over the DEFAULT-layout core -- residual_add 4096^3 +54.0%,
    bias_residual_gelu 8192^3 +50.6%, silu 4096^3 +9.5% (the source-C/residual family
    where the core caps ab at 2 and the FFI ring recovers it). The autotuner keeps
    whichever of {core, FFI alt} wins.

[A] No-autotune default -- what a user gets with autotuning OFF (default_config()),
    the promoted formula config vs the raw fragment default:
  * 64x8192x8192 fp8 321.9 vs 29.3 TF/s (~11x); 512x4096x4096 bf16 697.2 vs 557.5;
    2048x4096x4096 bf16 1340.5 vs 1289.8. This is the primary user-visible win of the
    seed and does not depend on running the search.

[B] Effect on the autotune SEARCH (formula seed ON vs OFF, both on this full stack):
  * Convergence: on 64x8192x8192 fp8 the seed lands the search AT the answer -- the
    first benchmarked config is ~1.05x the final best (seeded) vs ~4.77x (unseeded),
    reproducible across two reps. Both arms ultimately converge to the SAME winner
    ([64,64,128] cm1 ab12), re-benched cold-L2 to an identical 321.9 TF/s (any
    end-of-search delta in the tables is the autotuner's noisy in-process timer, not a
    real difference -- the cold-L2 re-bench is the trustworthy comparison).
  * Ceiling: given PRs 1-3, the unseeded search on this stack already REACHES the good
    cluster_m=2 configs (e.g. 512x4096x4096 bf16 base and treat both reach
    [256,128,128] cm2 ab4 = 697.2). That reachability is PR3's contribution (widening
    the search space), which the seed-OFF arm confirms -- NOT a win attributable to
    this seed. The formula's marginal search-ceiling effect is small (a soft +3.9% on
    2048x4096x4096 bf16, where the seeded search reached the deep-AB [256,256,64] ab6
    tile in 1 of 2 reps and the unseeded did not). The seed's real value is the
    no-autotune default and convergence above, not moving the search ceiling.

Prerequisite edits bundled here (needed for the formula to emit valid configs):
  * cute_mma.py: admit a rank-2 exact-shape residual aux on the explicit-epi-tile path
    (broadcast_axis in (1, None)) for the bias_residual_gelu / residual_add family --
    without it the FFI alt-seed's residual config is rejected (it drives the +50% wins).
  * tcgen05_config.py: extend the 16-bit deep-AB (ab>3) SMEM-clamped admission from
    cluster_m=2 to cluster_m=1 DEFAULT layout, so the formula's bf16 decode tile
    ([64,32,256] ab4 = 196608 B) validates. max_ab_stages_that_fit still rejects the
    real cm1 256^2 overflow (clamped to ab<=3).
  * _gelu_tanh_approx.py: lower cute.math.tanh with fastmath=True (hw tanh.approx.f32),
    recovering ~14% on the fused GEMM+GELU epilogue.

Also pins the promoted DEFAULT-layout seed's layout/ffi/epi-tile keys explicitly: on a
direct-entry-eligible shape the fragment defaults are FFI-biased, so a promoted seed
that omitted them merged into an invalid explicit_epi_tile+ffi hybrid and raised
InvalidConfig in the no-autotune / baseline compile path.

Known limitation, NOT fixed (out of scope): fp8 512x2048x2048 -- the formula emits
l2_groupings=[4] where the old producer emits [1]. Measured this run it is a near-tie
(seed/old 0.997), but the l2_groupings divergence is real; left for a follow-up.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

stack-info: PR: #3152, branch: calebmkim/stack/36
calebmkim added a commit that referenced this pull request Jul 24, 2026
Replaces the special-case cluster_m=2 seed producers with one shape-aware analytical
formula (cute analog of the Triton B200 formula heuristic #3007):
f(M,K,N,dtype,epilogue,num_sm,smem_budget) -> Config, promoted to the no-autotune
default. Covers configs the old producers structurally could not emit: cluster_m=1
decode (M<=128, previously NO seed) and deep-AB compute (ab=6 vs the old ab=3).

Depends on PRs 1-3 below it (they widen the search surface / unblock injection); this
PR is the config PRODUCER on top of that surface.

STATUS: this is a STARTING POINT, not a finished heuristic. It is a first, deliberately
simple analytical formula (~5 knobs) that already covers regimes the old producers
could not and beats them on quality (below), but it is expected to be iterated on:
the classifier thresholds (e.g. the wave-fill / _MANY_WAVE crossovers), the l2_groupings
rule, and coverage of more epilogues/shapes are all approximations with known gaps (see
the F1 limitation below). Intent is to land a maintainable baseline that replaces the
special-case producers and then improve it incrementally, not to ship a final tuning oracle.

What this PR is responsible for (measured this run, single B200, cold-L2 cudagraph):

[A] Config quality -- the config the formula EMITS vs the old producers / an external
    reference, benched directly (no autotuner):
  * geomean(seed / old-producer) = 1.091 over the 23 curriculum shapes where an old
    producer fired, 0 regressions (0 shapes >2% slower). The old producers emit NO
    config on the other 12 shapes (cluster_m=1 decode etc.) -- the "A1 gap" this
    formula closes.
  * geomean(seed / external ref) = 1.035 (curriculum), 1.083 (held-out fp16, 3 shapes
    with a ref) -- held-out >= curriculum, so no overfitting.
  * accuracy 35/35 PASS (curriculum core; 48/48 incl. the FFI alt-seed arm; 4/4 fp16).
  * FFI alt-seed KEPT: 3 curriculum shapes show a >2% win of the FFI explicit_epi_tile
    topology over the DEFAULT-layout core -- residual_add 4096^3 +54.0%,
    bias_residual_gelu 8192^3 +50.6%, silu 4096^3 +9.5% (the source-C/residual family
    where the core caps ab at 2 and the FFI ring recovers it). The autotuner keeps
    whichever of {core, FFI alt} wins.

[A] No-autotune default -- what a user gets with autotuning OFF (default_config()),
    the promoted formula config vs the raw fragment default:
  * 64x8192x8192 fp8 321.9 vs 29.3 TF/s (~11x); 512x4096x4096 bf16 697.2 vs 557.5;
    2048x4096x4096 bf16 1340.5 vs 1289.8. This is the primary user-visible win of the
    seed and does not depend on running the search.

[B] Effect on the autotune SEARCH (formula seed ON vs OFF, both on this full stack):
  * Convergence: on 64x8192x8192 fp8 the seed lands the search AT the answer -- the
    first benchmarked config is ~1.05x the final best (seeded) vs ~4.77x (unseeded),
    reproducible across two reps. Both arms ultimately converge to the SAME winner
    ([64,64,128] cm1 ab12), re-benched cold-L2 to an identical 321.9 TF/s (any
    end-of-search delta in the tables is the autotuner's noisy in-process timer, not a
    real difference -- the cold-L2 re-bench is the trustworthy comparison).
  * Ceiling: given PRs 1-3, the unseeded search on this stack already REACHES the good
    cluster_m=2 configs (e.g. 512x4096x4096 bf16 base and treat both reach
    [256,128,128] cm2 ab4 = 697.2). That reachability is PR3's contribution (widening
    the search space), which the seed-OFF arm confirms -- NOT a win attributable to
    this seed. The formula's marginal search-ceiling effect is small (a soft +3.9% on
    2048x4096x4096 bf16, where the seeded search reached the deep-AB [256,256,64] ab6
    tile in 1 of 2 reps and the unseeded did not). The seed's real value is the
    no-autotune default and convergence above, not moving the search ceiling.

Prerequisite edits bundled here (needed for the formula to emit valid configs):
  * cute_mma.py: admit a rank-2 exact-shape residual aux on the explicit-epi-tile path
    (broadcast_axis in (1, None)) for the bias_residual_gelu / residual_add family --
    without it the FFI alt-seed's residual config is rejected (it drives the +50% wins).
  * tcgen05_config.py: extend the 16-bit deep-AB (ab>3) SMEM-clamped admission from
    cluster_m=2 to cluster_m=1 DEFAULT layout, so the formula's bf16 decode tile
    ([64,32,256] ab4 = 196608 B) validates. max_ab_stages_that_fit still rejects the
    real cm1 256^2 overflow (clamped to ab<=3).
  * _gelu_tanh_approx.py: lower cute.math.tanh with fastmath=True (hw tanh.approx.f32),
    recovering ~14% on the fused GEMM+GELU epilogue.

Also pins the promoted DEFAULT-layout seed's layout/ffi/epi-tile keys explicitly: on a
direct-entry-eligible shape the fragment defaults are FFI-biased, so a promoted seed
that omitted them merged into an invalid explicit_epi_tile+ffi hybrid and raised
InvalidConfig in the no-autotune / baseline compile path.

Known limitation, NOT fixed (out of scope): fp8 512x2048x2048 -- the formula emits
l2_groupings=[4] where the old producer emits [1]. Measured this run it is a near-tie
(seed/old 0.997), but the l2_groupings divergence is real; left for a follow-up.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

stack-info: PR: #3152, branch: calebmkim/stack/36
@calebmkim
calebmkim force-pushed the calebmkim/stack/35 branch from 08071ba to b2dea58 Compare July 24, 2026 17:09
@calebmkim
calebmkim force-pushed the calebmkim/stack/36 branch from de6b9c2 to 844d1c4 Compare July 24, 2026 17:09
@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Meta Open Source bot. label Jul 24, 2026
@calebmkim
calebmkim changed the base branch from calebmkim/stack/35 to main July 27, 2026 22:13
calebmkim added a commit that referenced this pull request Jul 27, 2026
Replaces the special-case cluster_m=2 seed producers with one shape-aware analytical
formula (cute analog of the Triton B200 formula heuristic #3007):
f(M,K,N,dtype,epilogue,num_sm,smem_budget) -> Config, promoted to the no-autotune
default. Covers configs the old producers structurally could not emit: cluster_m=1
decode (M<=128, previously NO seed) and deep-AB compute (ab=6 vs the old ab=3).

Depends on PRs 1-3 below it (they widen the search surface / unblock injection); this
PR is the config PRODUCER on top of that surface.

STATUS: this is a STARTING POINT, not a finished heuristic. It is a first, deliberately
simple analytical formula (~5 knobs) that already covers regimes the old producers
could not and beats them on quality (below), but it is expected to be iterated on:
the classifier thresholds (e.g. the wave-fill / _MANY_WAVE crossovers), the l2_groupings
rule, and coverage of more epilogues/shapes are all approximations with known gaps (see
the F1 limitation below). Intent is to land a maintainable baseline that replaces the
special-case producers and then improve it incrementally, not to ship a final tuning oracle.

What this PR is responsible for (measured this run, single B200, cold-L2 cudagraph):

[A] Config quality -- the config the formula EMITS vs the old producers / an external
    reference, benched directly (no autotuner):
  * geomean(seed / old-producer) = 1.091 over the 23 curriculum shapes where an old
    producer fired, 0 regressions (0 shapes >2% slower). The old producers emit NO
    config on the other 12 shapes (cluster_m=1 decode etc.) -- the "A1 gap" this
    formula closes.
  * geomean(seed / external ref) = 1.035 (curriculum), 1.083 (held-out fp16, 3 shapes
    with a ref) -- held-out >= curriculum, so no overfitting.
  * accuracy 35/35 PASS (curriculum core; 48/48 incl. the FFI alt-seed arm; 4/4 fp16).
  * FFI alt-seed KEPT: 3 curriculum shapes show a >2% win of the FFI explicit_epi_tile
    topology over the DEFAULT-layout core -- residual_add 4096^3 +54.0%,
    bias_residual_gelu 8192^3 +50.6%, silu 4096^3 +9.5% (the source-C/residual family
    where the core caps ab at 2 and the FFI ring recovers it). The autotuner keeps
    whichever of {core, FFI alt} wins.

[A] No-autotune default -- what a user gets with autotuning OFF (default_config()),
    the promoted formula config vs the raw fragment default:
  * 64x8192x8192 fp8 321.9 vs 29.3 TF/s (~11x); 512x4096x4096 bf16 697.2 vs 557.5;
    2048x4096x4096 bf16 1340.5 vs 1289.8. This is the primary user-visible win of the
    seed and does not depend on running the search.

[B] Effect on the autotune SEARCH (formula seed ON vs OFF, both on this full stack):
  * Convergence: on 64x8192x8192 fp8 the seed lands the search AT the answer -- the
    first benchmarked config is ~1.05x the final best (seeded) vs ~4.77x (unseeded),
    reproducible across two reps. Both arms ultimately converge to the SAME winner
    ([64,64,128] cm1 ab12), re-benched cold-L2 to an identical 321.9 TF/s (any
    end-of-search delta in the tables is the autotuner's noisy in-process timer, not a
    real difference -- the cold-L2 re-bench is the trustworthy comparison).
  * Ceiling: given PRs 1-3, the unseeded search on this stack already REACHES the good
    cluster_m=2 configs (e.g. 512x4096x4096 bf16 base and treat both reach
    [256,128,128] cm2 ab4 = 697.2). That reachability is PR3's contribution (widening
    the search space), which the seed-OFF arm confirms -- NOT a win attributable to
    this seed. The formula's marginal search-ceiling effect is small (a soft +3.9% on
    2048x4096x4096 bf16, where the seeded search reached the deep-AB [256,256,64] ab6
    tile in 1 of 2 reps and the unseeded did not). The seed's real value is the
    no-autotune default and convergence above, not moving the search ceiling.

Prerequisite edits bundled here (needed for the formula to emit valid configs):
  * cute_mma.py: admit a rank-2 exact-shape residual aux on the explicit-epi-tile path
    (broadcast_axis in (1, None)) for the bias_residual_gelu / residual_add family --
    without it the FFI alt-seed's residual config is rejected (it drives the +50% wins).
  * tcgen05_config.py: extend the 16-bit deep-AB (ab>3) SMEM-clamped admission from
    cluster_m=2 to cluster_m=1 DEFAULT layout, so the formula's bf16 decode tile
    ([64,32,256] ab4 = 196608 B) validates. max_ab_stages_that_fit still rejects the
    real cm1 256^2 overflow (clamped to ab<=3).
  * _gelu_tanh_approx.py: lower cute.math.tanh with fastmath=True (hw tanh.approx.f32),
    recovering ~14% on the fused GEMM+GELU epilogue.

Also pins the promoted DEFAULT-layout seed's layout/ffi/epi-tile keys explicitly: on a
direct-entry-eligible shape the fragment defaults are FFI-biased, so a promoted seed
that omitted them merged into an invalid explicit_epi_tile+ffi hybrid and raised
InvalidConfig in the no-autotune / baseline compile path.

Known limitation, NOT fixed (out of scope): fp8 512x2048x2048 -- the formula emits
l2_groupings=[4] where the old producer emits [1]. Measured this run it is a near-tie
(seed/old 0.997), but the l2_groupings divergence is real; left for a follow-up.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

stack-info: PR: #3152, branch: calebmkim/stack/36
@calebmkim
calebmkim force-pushed the calebmkim/stack/36 branch 2 times, most recently from bcd9605 to 95ec8eb Compare July 27, 2026 22:14
@calebmkim
calebmkim changed the base branch from main to calebmkim/stack/35 July 27, 2026 22:14
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Meta Open Source bot.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant