[cute] Deep AB staging for fp8 to close the compute-bound gap

[yushangdi]

Stacked PRs:

• [cute] Unified rolled TMA producer + hoisted K-loop predicates for tcgen05 #2754
• [cute] Fused-scale epilogue: scalar colvec read for per-row scale #2742
• ->[cute] Deep AB staging for fp8 to close the compute-bound gap #2741
• [cute] Support column-major (K-major) B on the tcgen05 fp8 TMA path #2740

---

[cute] Deep AB staging for fp8 to close the compute-bound gap

With K-major B in place, the fp8 tcgen05 kernel was still capped at
ab_stages=3 (the bf16-tuned limit), so its software pipeline was too
shallow to hide K-loop TMA latency on compute-bound shapes (4096^3:
1143 TFLOP/s, 62% of torch._scaled_mm). 1-byte fp8 operands fit a much
deeper AB pipeline than 2-byte bf16 in the same SMEM budget.

Backports the tcgen05_config.py deep-staging logic from
#2696:

• max_ab_stages_that_fit(): largest ab_stages whose AB SMEM fits the
  per-CTA budget (mirrors CUTLASS _compute_stages).
• _validate_target1_ab_stage_envelope(): admit ab_stages>3 for fp8 as
  long as the AB SMEM fits, instead of hard-failing at 3.
• optional_fragments(): widen the ab_stages search cap to 12 for fp8
  (1-byte operands) on BOTH the search and validation surfaces, so the
  autotuner can sample deep pipelines and a frozen deep-staged fp8 config
  also passes normalize().

Benchmark (B200, CUDA 13.2, fp8 e4m3 scaled_mm, m=k=n=4096, col-major B,
CtaGroup.TWO cluster_m=2 role-local, do_bench, 10s warmup):

ab_stages= 3 (prev cap) : 1143 TFLOP/s 62% of aten
ab_stages= 6 : 1494 TFLOP/s 80%
ab_stages= 8 : 1613 TFLOP/s 86% <- sweet spot
ab_stages=10/12 : ~1600 TFLOP/s (SMEM-bound plateau)
torch._scaled_mm : 1867 TFLOP/s

So K-major B (prior commit) + deep AB staging together take fp8 scaled_mm
from ~31% to ~86% of torch._scaled_mm on the 4096^3 compute-bound shape.
Correctness rel_err 0.0000 throughout; full cute suite: 93 passed.

Co-Authored-By: Claude Fable 5 noreply@anthropic.com