[BUG][CuTe DSL] PipelineTmaUmma arrival masks include an out-of-cluster "v-peer" bit when the cluster's v-extent is 1, causing CGA out-of-range faults (Xid 13)

Open
  • Is something wrong?
#3,658 2 comments 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 4 days

@yunweili3 is already working on this.

Since Sep 22, 2026.

  • #3661 by @yunweili3 — open

Assessment

Difficulty
3/5
Estimated time
1-2 days
Newbie friendliness
40/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
cpp, python
Domain
backend, compilers

Research direction

The bug is in python/CuTeDSL/cutlass/pipeline/sm100.py, specifically the _compute_mcast_arrival_mask function. Start by understanding the cluster layout and how masks are computed for TMA multicast. The fix involves adding a condition to only include peer masks when the v-extent is 2, similar to PipelineAsyncUmma.create. Test with a kernel using cta_layout_vmnk=(1, 4, 1, 1) to ensure no CGA faults occur.

Written by the indexing model from the issue text.

Description

inactive-30d

Environment

  • nvidia-cutlass-dsl 4.6.2 (the code is unchanged at current main)
  • CUDA 13.0, Blackwell (sm_103a)

Summary

PipelineTmaUmma._compute_mcast_arrival_mask
(python/CuTeDSL/cutlass/pipeline/sm100.py, L101-L157)
unconditionally ORs in "peer" image masks computed at the coordinate

block_in_cluster_coord_vmnk_peer = (
    cta_in_cluster_coord_vmnk[0] ^ 1,
    *cta_in_cluster_coord_vmnk[1:],
)

The v-peer only exists for tcgen05.CtaGroup.TWO (v-extent 2). For a plain
non-paired cluster, e.g. cta_layout_vmnk = (1, C, 1, 1) — the natural
configuration for a TMA load multicast shared by C CTAs that each run their
own MMA — the peer coordinate (1, b, 0, 0) lies outside the layout's
domain (mode 0 has extent 1). With the compact strides (1, 1, C, C) it
evaluates to rank b + 1, so the mode-1 peer image is bits {1..C}. The
resulting producer/consumer mask is {0..C}: it contains bit C, a CTA
rank that does not exist in the cluster.

PipelineTmaUmma.create computes these masks whenever
cute.size(cta_layout_vmnk) > 1
(L290-L300)
with no check on the v-extent, and consumer_mask = producer_mask. The same
helper is used by PipelineTmaMultiConsumersAsync.create
(L1244),
which is exposed the same way.

By contrast, PipelineAsyncUmma.create implements the intended pattern: it
only computes peer masks when the pipeline actually uses 2-CTA UMMA
("No mcast mask if we're not using 2CTA tcgen05 MMA").

Observed behavior

A kernel with a TMA-load producer (data multicast along mode 1, mask built
with cpasync.create_tma_multicast_mask(..., mcast_mode=1)) and tcgen05
UMMA consumers, on a 4-CTA cluster created as
PipelineTmaUmma.create(..., cta_layout_vmnk=(1, 4, 1, 1)):

  • The kernel faults at cluster scope once consumers start releasing buffers.
    dmesg shows Xid 13, "SM Cga Exception: OOR Addr" — an SM-side
    cluster (mapa/remote-mbarrier) out-of-range access, matching an
    mbarrier.arrive aimed at rank 4 of a 4-CTA cluster.
  • An ablation that arrives with a self-only mask instead converts the crash
    into a hang (remote arrivals missing), confirming the mask — not the TMA
    transaction accounting — is at fault.

Workaround

Overriding both masks right after create() with the plain mode-1 image
mask fixes the fault; results are then bit-exact against the equivalent
non-clustered kernel:

coord = cta_layout_vmnk.get_flat_coord(cta_rank_in_cluster)
mask = cpasync.create_tma_multicast_mask(cta_layout_vmnk, coord, mcast_mode=1)
object.__setattr__(pipeline, "consumer_mask", mask)   # frozen dataclass
object.__setattr__(pipeline, "producer_mask", mask)

Suggested fix

Include the *_peer terms only when the v-extent is 2, e.g.:

if cute.size(cta_layout_vmnk, mode=[0], loc=loc, ip=ip) == 2:
    mask |= peer_terms

mirroring the guard that PipelineAsyncUmma.create already applies.

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.