[BUG][CuTeDSL] Blackwell dense_gemm example aborts (SIGABRT in cute.slice_) for tf32 + 2-CTA + B K-major with MMA tiler N 96 / 160, which can_implement() accepts

Open
  • Is something wrong?
#3,784 1 comment 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 4 days

@yunweili3 is already working on this.

Since Oct 25, 2026.

  • #3785 by @yunweili3 — open
  • #3791 by @yunweili3 — open

Assessment

Difficulty
4/5
Estimated time
3-5 days
Newbie friendliness
25/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
cpp, python
Domain
hpc

Research direction

Start from examples/python/CuTeDSL/cute/blackwell/kernel/dense_gemm/dense_gemm.py:398 where cute.slice_ aborts on b_smem_layout_staged. Run python repro.py --example dense_gemm.py with CUTE_DSL_ARCH=sm_100a to reproduce the SIGABRT. Check linked PR #3785 and #3468 for prior fixes; 'done' means the config either compiles, is rejected by can_implement(), or slice_ raises a clean error instead of aborting.

Written by the indexing model from the issue text.

Description

Describe the bug

The Blackwell dense_gemm.py CuTe DSL example aborts the Python process (SIGABRT, no Python exception, no MLIR
diagnostic) in cute.compile for tf32 operands with use_2cta_instrs=True, B K-major and an MMA tiler N of 96 or 160.
DenseGemmKernel.can_implement() accepts these configurations. The abort happens on the host, while the example builds
its IR, so no GPU is needed to see it.

The abort is in cute.slice_ at dense_gemm.py:398:

b_smem_layout = cute.slice_(self.b_smem_layout_staged, (None, None, None, 0))

The layout it slices (printed just before the call) is:

tiler (256, 96):  S<3,4,3> o 0 o ((48,8),1,4,8):((32,1),0,8,1536)   -> Aborted
tiler (256, 160): S<3,4,3> o 0 o ((80,8),1,4,7):((32,1),0,8,2560)   -> Aborted
tiler (256, 192): S<3,4,3> o 0 o ((96,8),1,4,6):((32,1),0,8,3072)   -> OK
tiler (256, 96), A operand: S<3,4,3> o 0 o ((128,8),1,4,8):((32,1),0,8,4096) -> OK

So it looks like the same SliceOp abort as #3454 (static swizzled ComposedLayout, first mode 10 aborts and 8 works),
but reached from NVIDIA's own example on a configuration that can_implement() accepts. #3454's repro also still aborts
on 4.8.1. The open PR #3468 ("Fixes #3454") would turn the abort into an unable to slice error. These example
configs would then still fail to compile, so either can_implement() or the B smem layout of the example needs a
change too.

Steps/Code to reproduce bug

repro.py (attached) loads examples/python/CuTeDSL/cute/blackwell/kernel/dense_gemm/dense_gemm.py, makes fake
tensors (cute.runtime.make_fake_compact_tensor), calls can_implement(), then cute.compile for sm_100a. Each
configuration runs in a child process.

pip install nvidia-cutlass-dsl==4.8.1
python repro.py --example dense_gemm.py            # tf32, 2-CTA, tiler (256, 96), cluster (2, 1), B K-major -> SIGABRT
python repro.py --example dense_gemm.py --sweep

Minimal form of what repro.py does:

import cutlass, cutlass.cute as cute, cuda.bindings.driver as cuda
from cutlass.cute.runtime import make_fake_compact_tensor
from dense_gemm import DenseGemmKernel          # CUTLASS examples/python/CuTeDSL/cute/blackwell/kernel/dense_gemm/

M, N, K, L = 512, 192, 128, 1
a = make_fake_compact_tensor(cutlass.TFloat32, (M, K, L), stride_order=(1, 0, 2), assumed_align=16)   # K-major
b = make_fake_compact_tensor(cutlass.TFloat32, (N, K, L), stride_order=(1, 0, 2), assumed_align=16)   # K-major
c = make_fake_compact_tensor(cutlass.Float16, (M, N, L), stride_order=(1, 0, 2), assumed_align=16)
gemm = DenseGemmKernel(cutlass.Float32, True, (256, 96), (2, 1), True)   # acc, use_2cta_instrs, mma_tiler_mn, cluster, TMA store
assert gemm.can_implement(a, b, c)
cute.compile(gemm, a, b, c, cuda.CUstream(0))  # Aborted (core dumped)

(run with CUTE_DSL_ARCH=sm_100a on a machine without a GPU)

Sweep with nvidia-cutlass-dsl 4.8.1 (tf32 A/B, f32 accumulator, cluster (2, 1), TMA store; all accepted by can_implement()):

MMA tiler M B major N 32 N 64 N 96 N 128 N 160 N 192 N 256
256 K ok ok abort ok abort ok ok
256 N ok ok ok ok ok ok ok
128 K ok ok abort ok abort ok ok
128 N ok ok ok ok ok ok ok

C dtype and C major do not matter (f16, bf16, f32, and fp8 M-major C abort at (256, 96)). We also ran a larger generated
sweep on 4.7.0: 113 tf32 2-CTA configs that can_implement() accepts (A major, C dtype / major, cluster (2,1) / (2,2) /
(4,1), TMA store on / off, M 128 / 256). The 18 configs with B K-major and N 96 or 160 abort, and again 18 of 18 abort
(rc 134) on 4.8.1. The other 95 compile: 93 with the IR dump on (CUTE_DSL_KEEP=ptx,ir,cubin), and 2 N-32 configs
with fp8 M-major C only without it (with it, the dump step fails to re-parse its own IR; repro.py --n 32 --c Float8E5M2 --c-major m compiles). 4.6.0 gives the same abort on the 4 configs we tried. repro.py removes CUTE_DSL_KEEP, so
the abort needs no debug option.

Expected behavior

Either the example compiles (PTX allows dense tf32 cta_group::2 MMA with N in steps of 16), or can_implement()
rejects these tiler shapes, or cute.slice_ raises a Python / MLIR error instead of aborting the process.

Environment details

  • nvidia-cutlass-dsl 4.8.1 (PyPI; also 4.7.0, 4.6.0), Python 3.12, Linux x86_64, CUTE_DSL_ARCH=sm_100a, CPU only (no GPU).
  • CUTLASS main 91843b71d90737e0a9b86f7d9df213cff5a3526f (2026-10-20); python/CuTeDSL/requirements.txt pins 4.8.1.

Additional context

Python stack (faulthandler) ends in _cute_ops_gen.py SliceOp.__init__ <- cute/core.py slice_ <-
dense_gemm.py:398 __call__ <- base_dsl/dsl.py build_ir_module. dense_gemm_software_pipeline.py:395 has the
same slice.

Dominant language
C++
Stars
10.7k
Forks
2.2k
Avg merge
3d 20h
Merged PRs (30d)
4

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

More from NVIDIA/cutlass

All issues in NVIDIA/cutlass

Similar issues

More C++ issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.