[CuTeDSL] make_tiled_mma SIGABRTs on hierarchical rank-3 atom_layout_mnk

Summary

cute.make_tiled_mma passes its Python-side rank check but crashes with SIGABRT inside the MLIR backend when given a rank-3 atom_layout_mnk whose M or N mode has a nested (hierarchical) shape. The equivalent layout works correctly in C++ CuTe.

Environment

• nvidia-cutlass-dsl (pip, latest)
• GPU: NVIDIA RTX 5080 (SM 120 / Blackwell)
• CUDA 13.1, Python 3.12, Linux

Minimal Repro

import cutlass
import cutlass.cute as cute
from cutlass.cute import nvgpu

@cute.jit
def test():
    mma_op = cute.make_mma_atom(nvgpu.MmaUniversalOp(cutlass.Float32))

    # Flat rank-3 layout — works fine
    flat = cute.make_layout((16, 8, 1), stride=(1, 16, 0))
    tiled_mma_flat = cute.make_tiled_mma(mma_op, flat)
    cute.printf("flat ok, size = {}", cute.size(tiled_mma_flat))

    # Hierarchical rank-3 layout — crashes with SIGABRT
    hier = cute.make_layout(((8, 2), (4, 2), 1),
                             stride=((1, 32), (8, 64), 0))
    tiled_mma_hier = cute.make_tiled_mma(mma_op, hier)
    cute.printf("hier ok, size = {}", cute.size(tiled_mma_hier))

test()

Expected

Both make_tiled_mma calls should succeed. This layout is the canonical way to express warp-structured thread tiling in CuTe C++ (e.g., warptiling SGEMM where ((NTM, NWM), (NTN, NWN)) encodes NWM*NWN warps each owning an NTM*NTN = 32 thread subtile).

Actual

python repro.py
# Aborted (core dumped), exit=134

No Python exception — the process dies inside MLIR type construction. Neither stderr nor any Python try/except catches it.

Additional context

• cute.make_layout(((8,2),(4,2)), stride=((1,32),(8,64))) (rank-2 hierarchical) constructs fine outside make_tiled_mma, so the layout itself is valid.
• make_tiled_copy_tv accepts similar hierarchical shapes without crashing.
• C++ CuTe accepts the exact same shape/stride via `make_layout(make_shape(make_shape(_8,_2), make_shape(_4,_2), _1), make_stride(...))` and make_tiled_mma produces a working warp-structured MMA.
• permutation_mnk is not a workaround — it reshuffles atom value positions but cannot express warp groupings in the thread mapping (tested).

The DSL-side check at cutlass/cute/atom.py:544 only validates rank(atom_layout_mnk) != 3, so hierarchical layouts pass through to MLIR which then asserts.

Start at cutlass/cute/atom.py:544 and reproduce with the provided Python script, then inspect the MLIR type construction reached by hierarchical rank-3 layouts. Compare the working flat and rank-2 cases with make_tiled_copy_tv; done means both make_tiled_mma calls complete without SIGABRT and produce the expected tiled MMA.