[BUG] Grouped GEMM and grouped Rank2K kernels build operand iterators from a temporary Params

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

Maintainers usually reply within 4 days

@BwL1289 is already working on this.

Since Oct 7, 2026.

  • #3719 by @BwL1289 — open

Assessment

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

Research direction

Read gemm/kernel/gemm_grouped.h, gemm_grouped_per_group_scale.h, gemm_grouped_softmax_mainloop_fusion.h, and rank_2k_grouped.h, focusing on operand iterator construction and Params lifetimes. Reproduce the reported clang -O2 builds for examples 24 and 38, then make the operand Params outlive each iterator. Done means the grouped GEMM and Rank2K PTX no longer show the reported unwritten-register addresses.

Written by the indexing model from the issue text.

Description

CUTLASS C++

gemm/kernel/gemm_grouped.h builds the mainloop operand iterators from a temporary Params:

typename Mma::IteratorA iterator_A(
  LayoutA(ldm_A),
  ptr_A, {problem_size.m(), problem_size.k()}, thread_idx, tb_offset_A);

LayoutA(ldm_A) becomes a temporary IteratorA::Params that is destroyed at the end of the statement, but the pitch-linear PredicatedTileAccessIterator keeps Params const &params_, so every later ++ / add_tile_offset reads a dead object. Same pattern in gemm_grouped_per_group_scale.h, gemm_grouped_softmax_mainloop_fusion.h, and rank_2k_grouped.h (spelled typename Mma1::IteratorA::Params(ldm_A) there). The epilogue iterators in rank_2k_grouped.h copy their Params, so those are fine.

nvcc's output happens not to show it. clang's (LLVM main, 24.0.0git) at -O2 does: building example 24's GemmGrouped kernel (half/half/float, Sm80 multistage, 128x128x32, 4 stages) for sm_121a, 30 of the 32 cp.async copies in the PTX take their global address from a register that is never written. Example 38's Rank2KGrouped kernel: 56 of 64.

Naming the Params so it lives as long as the iterator fixes it (0 of 32 and 0 of 64 after):

typename Mma::IteratorA::Params params_A{LayoutA(ldm_A)};  // braces: with parens this declares a function
typename Mma::IteratorA iterator_A(
  params_A, ptr_A, {problem_size.m(), problem_size.k()}, thread_idx, tb_offset_A);

Checked against main at 0b55a2f691 with CUDA 13.4. To get main through clang at all I had to pass -D__nv_atomic_load_n=__atomic_load_n -D__NV_ATOMIC_RELAXED=__ATOMIC_RELAXED (subbyte_reference.h uses the nvcc-only builtin); that's a separate problem.

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.