Skip to content

fmha_sm100(cute): do not depend on quack.activation.sub_packed_f32x2 - #27

Open
Orien-Xu wants to merge 1 commit into
MiniMax-AI:mainfrom
Orien-Xu:quack-sub-packed
Open

Orien-Xu wants to merge 1 commit into
MiniMax-AI:mainfrom
Orien-Xu:quack-sub-packed

Conversation

@Orien-Xu

@Orien-Xu Orien-Xu commented Oct 24, 2026 •

Copy link
Copy Markdown

Hi, I'm Orien from Parasail. This fixes an issue we ran into while serving MiniMax-M3.1-Flash (#24).

Fixes #24.

quack-kernels 0.6.5, which quack-kernels>=0.2.10 allows, removed
quack.activation.sub_packed_f32x2. ex2_emulation_2 in cute/src/common/utils.py still calls
it, so the CuTe kernels fail to compile with AttributeError.

Change

cute/src/common/utils.py: bind sub_packed_f32x2 to cute.arch.sub_packed_f32x2 when CUTLASS
DSL has it (the function quack 0.6.5 itself switched to), otherwise import the old quack helper.
The two call sites in ex2_emulation_2 use the local name. The arguments (two FP32 pairs, default
round-to-nearest-even) are unchanged. 10 lines changed, no other files.

Testing

  • TODO before submission (needs SM100): run the fmha_sm100 CuTe tests with
    (a) quack-kernels 0.6.4 + nvidia-cutlass-dsl 4.6.2 and (b) quack-kernels 0.6.5 +
    nvidia-cutlass-dsl >= 4.7, and check the outputs are bit-identical to the current main under (a).

GPU validation (NVIDIA B200, driver 595.71.05, CUDA 13.3, torch 2.14.1+cu130, head e9a19fe, base main 80434d7)

One causal sparse_atten_func call (Sq=Skv=4096, Hkv=1, qhead_per_kv=16, topK=16, BF16; causal enables ex2_emulation_2). CuTe compile caches cleared before every run.

quack-kernels / nvidia-cutlass-dsl main 80434d7 this PR
0.5.3 / 4.6.0.dev0 (cu12) TypeError: fmax() takes 2 positional arguments … OK, digest dce57afcfb852cd9
0.6.4 / 4.6.2 (cu12) TypeError: fmax() … OK, dce57afcfb852cd9
0.6.5 / 4.8.1 (cu12) TypeError: fmax() … OK, dce57afcfb852cd9
0.6.4 / 4.6.2 [cu13] OK, dce57afcfb852cd9 OK, dce57afcfb852cd9
0.6.5 / 4.8.1 [cu13] AttributeError: module 'quack.activation' has no attribute 'sub_packed_f32x2' OK, dce57afcfb852cd9

Output bytes are identical to main in the one configuration where main works.

pytest test_sparse_atten.py (the 60 BF16 test_sparse_atten cases, lse_temperature_scale=None, check_empty=False), quack 0.6.4 + cutlass-dsl[cu13] 4.6.2: main 24 passed / 36 failed, this PR 24 passed / 36 failed, identical failing set. The 36 failures already happen on main: all qhead_per_kv in {1, 2, 4}, an NVVM backend compilation error for sm_100a with CUDA 13.3. The PR passes the same 24 cases under quack 0.6.5 + cutlass-dsl[cu13] 4.8.1, where main cannot run at all.

…eases

Two incompatibilities stop the CuTe kernels from compiling with the newest
packages that requirements.txt allows (nvidia-cutlass-dsl>=4.4.1,
quack-kernels>=0.2.10):

- quack-kernels 0.6.5 removed quack.activation.sub_packed_f32x2 (it now uses
  cute.arch.sub_packed_f32x2 from CUTLASS DSL). ex2_emulation_2 still called
  it: AttributeError: module 'quack.activation' has no attribute
  'sub_packed_f32x2'.
- nvidia-cutlass-dsl >= 4.6.0 binds nvvm.fmax(a, b, *, c) with an inferred
  result type for every CUDA flavor, but fmax() chose the old
  nvvm.fmax(res, a, b) call whenever CUDA_VERSION is 12.9, which is what the
  default (cu12) install reports: TypeError: fmax() takes 2 positional
  arguments but 3 positional arguments ... were given.

Use cute.arch.sub_packed_f32x2 when CUTLASS DSL provides it and fall back to
the quack helper otherwise. Pick the nvvm.fmax calling convention from the
binding's signature instead of the CUDA version (MiniMax-AI#15 made the same fmax
change together with exact version pins).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@Orien-Xu
Orien-Xu marked this pull request as ready for review October 24, 2026 00:47
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

fmha_sm100 (CuTe path) fails with quack-kernels 0.6.5: quack.activation has no sub_packed_f32x2

1 participant