Skip to content

[autotuner][cute] repair partial FFI candidates away from the seed, not onto it - #3167

Draft
calebmkim wants to merge 1 commit into
calebmkim/stack/33from
calebmkim/stack/37
Draft

[autotuner][cute] repair partial FFI candidates away from the seed, not onto it#3167
calebmkim wants to merge 1 commit into
calebmkim/stack/33from
calebmkim/stack/37

Conversation

@calebmkim

@calebmkim calebmkim commented Jul 27, 2026

Copy link
Copy Markdown
Contributor

Stacked PRs:


[autotuner][cute] repair partial FFI candidates away from the seed, not onto it

The parent commit made the search draw tcgen05_tvm_ffi_launch False-only, but that
flag is just ONE of the six clauses in target1_tvm_ffi_promotion_requested. The
other five -- tcgen05_flat_role_coordinates, tcgen05_layout_strategy='explicit_epi_tile',
and the three tcgen05_layout_overrides
{epi_tile_m,epi_tile_n,d_store_box_n} -- are
independent search draws, so ~96% of candidates still tripped the gate and were
overwritten wholesale with the FFI seed (measured, S3 2048x4096x4096 bf16: 293/300
candidates -> 1 distinct config, 0 of 30 keys varying). That is not a search.

The projection was never needed for FFI coverage: the envelope is already in the
population as an eligibility-gated seed (CuteTcgen05ClusterM2FfiHeuristic plus the
formula FFI-alt), which _generate_best_available_population_flat additionally records
as a pinned finalist, so it is benchmarked and protected from pruning regardless.
Verified on 4 FFI-eligible shapes with the projection bypassed: the seed still reaches
the population every time, canonical-hash identical to the normalized seed.

_fix_target1_tvm_ffi_search_config becomes a three-way gate:

  1. A config already carrying the COMPLETE envelope (_is_coherent_target1_tvm_ffi_config
    over the 9 keys in _TARGET1_TVM_FFI_ENVELOPE_KEYS) is left untouched -- it is a
    valid direct-entry kernel, and projecting would clobber out-of-envelope knobs it
    legitimately chose. Not reachable from a random draw (0/2000 pure draws, since
    tvm_ffi_launch is drawn False-only and is an envelope key); this branch exists for
    the seed and for user-supplied / cache-restored configs.
  2. An EXPLICIT but partial tvm_ffi_launch=True (user config or cache entry; the search
    never draws True) is still projected onto the validated envelope -- unchanged
    behavior. An explicit request must not be silently downgraded, and a partial
    envelope cannot compile on its own.
  3. Everything else -- a search candidate that tripped the gate via one of the five
    other keys -- is repaired AWAY from FFI via _strip_target1_tvm_ffi_promotion_surface,
    preserving the ~20 other keys it sampled. The four layout keys move as one unit
    because validate_tcgen05_strategy_invariants couples them.

fix_search_config now runs the FFI stage FIRST, ahead of fix_cluster_m2_search_config.
Its bn-snapping guard only fires when the layout is already DEFAULT, so under the old
order a repaired candidate had already been forced to bn=256 and could never reach the
un-seeded [256,128,] tile. Measured contribution, same seeded population: [256,128,]
71 -> 139 (S3) and 72 -> 140 (medM). The original position (added by #2518) predates any
such dependency -- the old projection wiped every tcgen05
* key, so ordering was moot.

Effect (single B200, 300-draw populations, interleaved old/new control in one process):

shape distinct ffi=True cm2 [256,128,*]
S3 2048x4096x4096 bf16 18 -> 296 283 -> 1 7 -> 132
medM 512x4096x4096 bf16 18 -> 296 283 -> 1 8 -> 133
S3resid 2048x4096x4096 resid_add 19 -> 300 282 -> 1 9 -> 144
8192x1024x1024 bf16 11 -> 299 290 -> 1 4 -> 149

Correctness: 36 distinct post-repair configs compiled, run, and compared against
torch.matmul (atol=rtol=2e-2) -- 36/36 correct, zero wrong-number results. Repaired
candidates skip codegen slightly more often than naturally-non-FFI ones (20% vs ~13%,
all the same pre-existing non-whitelisted-fused-epilogue gate), which the autotuner
handles by skipping; net usable configs still rise by more than an order of magnitude.
Branch 1 verified directly: a coherent config with a mutated out-of-envelope knob
(l2_swizzle_size, l2_groupings, acc_stages, num_stages) keeps that knob and compiles
correct, while breaking one envelope key sends it back through branch 2.

Known adjacent issue, pre-existing and NOT addressed here: on residual_add shapes the
normalized FFI seed does not compile (flatten/unflatten rewrites ab_stages 3->2 and
c_stages 2->4 via _fix_aux_tma_full_tile_search_config, after which the flat-role
codegen guard rejects it). Confirmed identical on the unpatched parent.

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

…ot onto it

The parent commit made the search draw tcgen05_tvm_ffi_launch False-only, but that
flag is just ONE of the six clauses in _target1_tvm_ffi_promotion_requested. The
other five -- tcgen05_flat_role_coordinates, tcgen05_layout_strategy='explicit_epi_tile',
and the three tcgen05_layout_overrides_{epi_tile_m,epi_tile_n,d_store_box_n} -- are
independent search draws, so ~96% of candidates still tripped the gate and were
overwritten wholesale with the FFI seed (measured, S3 2048x4096x4096 bf16: 293/300
candidates -> 1 distinct config, 0 of 30 keys varying). That is not a search.

The projection was never needed for FFI coverage: the envelope is already in the
population as an eligibility-gated seed (CuteTcgen05ClusterM2FfiHeuristic plus the
formula FFI-alt), which _generate_best_available_population_flat additionally records
as a pinned finalist, so it is benchmarked and protected from pruning regardless.
Verified on 4 FFI-eligible shapes with the projection bypassed: the seed still reaches
the population every time, canonical-hash identical to the normalized seed.

_fix_target1_tvm_ffi_search_config becomes a three-way gate:
  1. A config already carrying the COMPLETE envelope (_is_coherent_target1_tvm_ffi_config
     over the 9 keys in _TARGET1_TVM_FFI_ENVELOPE_KEYS) is left untouched -- it is a
     valid direct-entry kernel, and projecting would clobber out-of-envelope knobs it
     legitimately chose. Not reachable from a random draw (0/2000 pure draws, since
     tvm_ffi_launch is drawn False-only and is an envelope key); this branch exists for
     the seed and for user-supplied / cache-restored configs.
  2. An EXPLICIT but partial tvm_ffi_launch=True (user config or cache entry; the search
     never draws True) is still projected onto the validated envelope -- unchanged
     behavior. An explicit request must not be silently downgraded, and a partial
     envelope cannot compile on its own.
  3. Everything else -- a search candidate that tripped the gate via one of the five
     other keys -- is repaired AWAY from FFI via _strip_target1_tvm_ffi_promotion_surface,
     preserving the ~20 other keys it sampled. The four layout keys move as one unit
     because validate_tcgen05_strategy_invariants couples them.

fix_search_config now runs the FFI stage FIRST, ahead of _fix_cluster_m2_search_config.
Its bn-snapping guard only fires when the layout is already DEFAULT, so under the old
order a repaired candidate had already been forced to bn=256 and could never reach the
un-seeded [256,128,*] tile. Measured contribution, same seeded population: [256,128,*]
71 -> 139 (S3) and 72 -> 140 (medM). The original position (added by #2518) predates any
such dependency -- the old projection wiped every tcgen05_* key, so ordering was moot.

Effect (single B200, 300-draw populations, interleaved old/new control in one process):

  shape                              distinct    ffi=True   cm2 [256,128,*]
  S3      2048x4096x4096 bf16       18 -> 296   283 -> 1      7 -> 132
  medM    512x4096x4096  bf16       18 -> 296   283 -> 1      8 -> 133
  S3resid 2048x4096x4096 resid_add  19 -> 300   282 -> 1      9 -> 144
          8192x1024x1024 bf16       11 -> 299   290 -> 1      4 -> 149

Correctness: 36 distinct post-repair configs compiled, run, and compared against
torch.matmul (atol=rtol=2e-2) -- 36/36 correct, zero wrong-number results. Repaired
candidates skip codegen slightly more often than naturally-non-FFI ones (20% vs ~13%,
all the same pre-existing non-whitelisted-fused-epilogue gate), which the autotuner
handles by skipping; net usable configs still rise by more than an order of magnitude.
Branch 1 verified directly: a coherent config with a mutated out-of-envelope knob
(l2_swizzle_size, l2_groupings, acc_stages, num_stages) keeps that knob and compiles
correct, while breaking one envelope key sends it back through branch 2.

Known adjacent issue, pre-existing and NOT addressed here: on residual_add shapes the
normalized FFI seed does not compile (flatten/unflatten rewrites ab_stages 3->2 and
c_stages 2->4 via _fix_aux_tma_full_tile_search_config, after which the flat-role
codegen guard rejects it). Confirmed identical on the unpatched parent.

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

stack-info: PR: #3167, branch: calebmkim/stack/37
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