[BUG] NVCC 13.3 cudafe++ explodes host code for structured bindings of CuTe TMA tuples

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug

With CUDA NVCC 13.3.73, using a structured binding to unpack a tuple of CuTe TMA copy objects and then using the binding names in kernel template arguments causes cudafe++ to generate an unexpectedly large host-side .cudafe1.cpp file.

The same source compiles successfully with NVCC 13.0.88. With NVCC 13.3.73, replacing the structured binding with an ordinary tuple variable plus std::get<N> also compiles successfully.

This reproduces with current CUTLASS main at 147295a3d4b75f3aeff247c25b8927cea9a7006a (version header reports 4.8.0), so it does not appear to be fixed by a newer CuTe header revision. The evidence points to a CUDA 13.3 cudafe++ regression, but I am reporting it here because the reproducer uses public CuTe TMA types.

Steps/Code to reproduce bug

#include <tuple>
#include <cute/tensor.hpp>

using namespace cute;

template <class Tensor>
auto make_three_tmas(Tensor tensor) {
  using T = bfloat16_t;
  using KAtom = GMMA::Layout_K_SW128_Atom<T>;
  using MNAtom = GMMA::Layout_MN_SW128_Atom<T>;
  using QLayout = decltype(tile_to_shape(KAtom{}, make_shape(_128{}, _192{})));
  using KVLayout = decltype(tile_to_shape(KAtom{}, make_shape(_128{}, _192{}, _2{})));
  using YLayout = decltype(tile_to_shape(MNAtom{}, make_shape(_64{}, _128{})));

  auto q = make_tma_copy(SM90_TMA_LOAD{}, tensor, QLayout{});
  auto kv = make_tma_copy(SM90_TMA_LOAD{}, tensor, take<0, 2>(KVLayout{}));
  auto y = make_tma_copy(SM90_TMA_STORE{}, tensor, YLayout{});
  return std::make_tuple(q, kv, y);
}

template <class Q, class KV, class Y>
__global__ void probe(Q, KV, Y) {}

void reproduce(void* ptr, int rows, int cols, int heads, int ld) {
  using T = bfloat16_t;
  auto tensor = make_tensor(make_gmem_ptr(static_cast<T*>(ptr)),
                            make_shape(rows, cols, heads),
                            make_stride(ld, _1{}, cols));
#ifdef WORKAROUND
  auto tmas = make_three_tmas(tensor);
  auto q = std::get<0>(tmas);
  auto kv = std::get<1>(tmas);
  auto y = std::get<2>(tmas);
#else
  auto [q, kv, y] = make_three_tmas(tensor);
#endif
  auto kernel = probe<decltype(q), decltype(kv), decltype(y)>;
  kernel<<<1, 1>>>(q, kv, y);
}

Compile from a CUTLASS checkout:

nvcc -std=c++20 -O3 \
  -gencode arch=compute_90a,code=sm_90a \
  --expt-relaxed-constexpr \
  -DCUTE_SM90_EXTENDED_MMA_SHAPES_ENABLED \
  -Iinclude --keep --keep-dir keep \
  -c repro.cu -o repro.o

The workaround build adds -DWORKAROUND.

Measured results with the public reproducer:

The 13.3 failing process was intentionally terminated once its generated host file exceeded 64 MiB to avoid unbounded disk usage. In the original larger translation unit, it produced a roughly 29 GB .cudafe1.cpp file before subsequent host compiler parse errors.

Expected behavior

NVCC 13.3 should compile the structured-binding form as NVCC 13.0 does, without expanding the bound CuTe types repeatedly into a huge generated host source file.

Environment details

• Environment location: bare metal
• OS: TencentOS Server 4, Linux x86_64
• Host compiler: GCC 12.3.1
• Failing compiler: CUDA NVCC 13.3.73
• Working comparison compiler: CUDA NVCC 13.0.88
• CUTLASS: current main, commit 147295a3d4b75f3aeff247c25b8927cea9a7006a, version 4.8.0
• C++ standard: C++20
• Target: sm_90a

Additional context

The original failure was first observed with GCC 14.3, but it also reproduces with GCC 12.3.1. Keeping the tuple object named and extracting its elements with std::get<N> avoids the CUDA 13.3 front-end expansion while preserving the types and generated kernel behavior.

If this is best tracked directly by the CUDA compiler team, guidance or internal routing to the appropriate NVCC owners would be appreciated.

Start with the public reproducer and compile command in the issue, comparing NVCC 13.3.73 structured bindings against the std::get workaround and NVCC 13.0.88. Confirm whether the behavior is isolated to cudafe++ rather than CUTLASS, then determine whether a CUTLASS change is possible or route the report to the CUDA compiler team.