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

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 0b55a2f 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.