develop: 39 new Blackwell (sm100) OSS test failures from #427 and #456 — both need a newer nvidia-cutlass-dsl

[Anerudhan]

Summary

develop regressed from 4 to 43 failing tests in the Blackwell (sm100) OSS test suite. Bisecting the five commits merged since the last good run attributes all 39 new failures to exactly two commits:

Commit	PR	New failures
b950af1	#427 — CSA: add fused Compressor forward+backward CuTe-DSL kernels	+14
57dc579	#408 — Relicense to Apache-2.0	0
5235c2b	#453 — BSA: fix cute-dsl bulk-copy elect gate	0
ad8f8b5	#457 — Tune Rubin FP8 K tiling and output staging	0
708ff87	#456 — Integrate Rubin wgrad cutedsl kernel	+25

Baseline for the bisect is 20d62d5 (#435), the last commit with a known-good run (2 failures per shard).

Both regressions have the same underlying cause: the new code calls CuTe-DSL APIs that do not exist in the nvidia-cutlass-dsl version the project's own dependency spec resolves to.

1. b950af1 (#427, CSA compressor) — 14 failures

Every one of the 14 fails identically:

E   TypeError: exp2() got an unexpected keyword argument 'approx'
python/cudnn/csa/compressor/compressor_sm100_r128.py:130: in _exp_fast
    return cute_math.exp2(y, approx=True, ftz=True)

Affected: all of test/python/fe_api/csa/test_CSA_compressor.py —
test_numerics_vs_references (7 params), test_backward_fills_uninitialized_buffers (4 params),
test_backward_grad_ape_zeroing_contract, test_cuda_graph_capture_r128, test_replay_determinism_r128.

cutlass.cute.math.exp2 has signature (a, fastmath=False, *, loc, ip) — there is no approx or ftz parameter.

2. 708ff87 (#456, Rubin wgrad) — 25 failures

24 × ICE during cute.compile:

cutlass.base_dsl.common.DSLRuntimeError: ICE IR Verification Failed
error: 'cute.get_shape' op using value defined outside the region
  at python/cudnn/gemm/.../moe_blockscaled_grouped_gemm_wgrad.py:638
     c_smem_layout_staged = sm100_utils.make_smem_layout_epi(...)
  note: "BlockScaledMoEGroupedGemmWgradKernel" (...:598): required by region isolation constraints

All 24 are the discrete-output wgrad paths:
test_grouped_gemm_wgrad_discrete_compile_execute_{fp4,fp8},
test_grouped_gemm_wgrad_discrete_accumulate_compile_execute_{fp4,fp8},
test_grouped_gemm_wgrad_discrete_wrapper_{fp4,fp8},
test_grouped_gemm_wgrad_dynamic_tokens_compile_execute_{fp4,fp8}[...-discrete]
(each × 3 cluster/tiler combinations).

1 × missing module:

test_rubin_kernel_dispatch.py::test_grouped_gemm_wgrad_rubin_tmem_plan_rejects_invalid_sf_vector
  ModuleNotFoundError: No module named 'cutlass.utils.rubin_helpers'

These same 41 tests pass at ad8f8b5 (the immediately preceding commit) and fail at 708ff87.

Root cause: unpinned cutlass-dsl floor

pyproject.toml specifies only a lower bound:

cutedsl = [ "nvidia-cutlass-dsl[cu13]>=4.5.0", ... ]

Neither 4.5.0 nor 4.5.2 provides cutlass.utils.rubin_helpers, and neither accepts approx= on cute.math.exp2 — verified directly:

cutlass-dsl 4.5.2:  rubin_helpers -> ModuleNotFoundError    exp2 accepts approx= -> False
cutlass-dsl 4.5.0:  rubin_helpers -> ModuleNotFoundError    exp2 accepts approx= -> False

So develop currently requires a cutlass-dsl newer than its own spec guarantees, and any fresh environment resolving the floor is broken. (Note there are separate oss:rel:cutlass_dsl_4.6 / cutlass_dsl_internal CI variants; the default oss:rel job is the one that is red.)

Suggested fix

Raise the floor (and ideally add an upper bound) on nvidia-cutlass-dsl in pyproject.toml to the minimum version that actually provides rubin_helpers and exp2(approx=), or make the two call sites version-tolerant.

Not caused by the GEMM reorg

For the record, PR #459 (GEMM fusion reorg) is a pure file move and is not implicated: the two failing kernel files are byte-identical across that PR, python/cudnn/csa/ is untouched by it, and the failures reproduce on 708ff87 with no reorg applied.

Pre-existing (not from these commits)

test_DSA_indexer_top_k (2 tests, test_..._compile_execute and test_..._wrapper) fails at every commit including the baseline, with a numerical mismatch (Mismatched elements: 422/705, Greatest absolute difference: inf). Tracked separately from this regression.

Reproduction

Blackwell sm100, cuDNN 9.25, one shared build with only the Python tree swapped per commit so the tree is the sole variable:

pytest test/python/fe_api/csa/test_CSA_compressor.py \
       test/python/fe_api/grouped_gemm/test_grouped_gemm_wgrad.py \
       test/python/fe_api/test_rubin_kernel_dispatch.py -m L0

commit	wgrad+rubin	CSA
20d62d5 … ad8f8b5	41 passed, 0 failed	14 failed from b950af1 onward
708ff87	25 failed, 30 passed	14 failed

cc @ the authors of #427 and #456.

[Anerudhan]

@zkyue and @WanZzzzzz Do you want to look at this?

@zkyue I think the requirement for the kernels is 4.6.0 and above cutedsl version. I would add that to the test and the code.

CSA tests: @zkyue
4.5.0 (CI pins this) 14 failed, 74 passed
4.6.1 88 passed

wgrad test: @WanZzzzzz
4.5.0 25 failed, 30 passed
4.6.1 25 failed, 30 passed

[zkyue]

Thanks for the bisect — the 14 CSA failures are ours (#427), and the attribution is right: the only other commit touching that file on develop is #408's license header.

Fixed at the call site rather than by raising the floor: #463. approx=/ftz= are newer than the >=4.5.0 floor, but fastmath=True is accepted by both versions and lowers to the same instruction, which is the part that matters here — the r128 tolerance contract is calibrated on ex2.approx.ftz.f32:

-    return cute_math.exp2(y, approx=True, ftz=True)
+    return cute_math.exp2(y, fastmath=True)

Verified on a B200 (CC 10.0, cuDNN 9.23) with cutlass-dsl 4.5.0 and 4.6.1 installed side by side:

• 4.5.0, before: test_CSA_compressor.py -m L0 → 14 failed, 74 passed (matches your table exactly).
• 4.5.0, after: 88 passed, 1 skipped.
• 4.6.1, after: the 16-kernel reg_probe_csa_compressor_r128.py PTX is byte-identical to before the change (diff -rq clean) — same register counts, 0 spill / 0 stack, 1104 ex2.approx.ftz.f32 across the set. Nothing changes for anyone already on 4.6.
• 4.5.0, after: the same 1104 ex2.approx.ftz.f32, with no no-.ftz ex2.approx.f32 anywhere — so the instruction the contract is calibrated on is unchanged at the floor too.
• gate_csa_compressor_r128.py (the script the tolerance contract is gated by): GATE PASS 21/21 on both 4.5.0 and 4.6.1 — padded/ragged shapes, the 131072-token cases, the NaN-overflow contract, and schedule coverage fwd 10/10 bwd 6/6.

No test skip or version gate is needed: the suite is green at the dependency floor.

On raising the floor itself — worth noting it wouldn't have cleared #456: those 25 wgrad failures reproduce on 4.6.1 in your own table, and cutlass.utils.rubin_helpers is a separate gap, so that half may be cleaner to track independently of this one.

[WanZzzzzz]

The rubin wgrad changes broke the GB300 test. Below are the fixes:

• Fixed Blackwell region-isolation ICE for discrete WGRAD paths by constructing TMA metadata outside the helper kernel on SM100.
• Preserved Rubin’s in-helper TMA construction.
• Fixed FP4 numerical mismatches with internal CUTLASS by bypassing the legacy FP4 layout recast; CUTLASS 4.5 retains it.
• Skipped Rubin-only dispatch validation when CUTLASS 4.5 lacks rubin_helpers.
• Updated imports and tests for the new cudnn.gemm.cutedsl.grouped.wgrad structure.

GB300 results:

Before:
CUTLASS 4.5: 25 failed, 30 passed
Internal CUTLASS: 33 failed, 22 passed
After:
CUTLASS 4.5: 48 passed, 15 skipped
Internal CUTLASS: 49 passed, 14 skipped

The 14 common skips are unsupported parameter combinations.
Test count increased from 55 to 63 due to the newer rebased develop suite.

[Anerudhan]

Both commits are merged.

Thanks closing