[BUG] `CUDA_CTA_RECONFIG_ACTIVATED` is not enabled for SM110a (Jetson AGX Thor)

Which component has the problem?

CuTe DSL

Bug Report

Problem

On Jetson AGX Thor (sm_110a), CTA register reconfiguration does not appear to be enabled by CUTLASS's architecture feature gate in include/cutlass/arch/reg_reconfig.h.

The current implementation enables CUDA_CTA_RECONFIG_ACTIVATED for several architectures, including SM90, SM100/101/103, SM107, and SM120/121, but there does not appear to be a corresponding SM110/SM110a case.

This becomes visible in downstream warp-specialized kernels that use setmaxnreg.inc / setmaxnreg.dec for producer-consumer register redistribution.

Environment

• Device: NVIDIA Jetson AGX Thor
• Target architecture: sm_110a
• CUDA compilation target: -arch=sm_110a
• Downstream framework: TileLang
• Workload: warp-specialized GEMM with software pipelining

Observed Behavior

A TileLang-generated warp-specialized GEMM uses CTA register redistribution between producer and consumer warp groups, for example:

tl::warpgroup_reg_dealloc<24>();
...
tl::warpgroup_reg_alloc<240>();

When compiling the generated CUDA code for Thor, compilation fails with:

tl::warpgroup_reg_dealloc requires a target with CTA register reconfiguration, such as sm_90a

tl::warpgroup_reg_alloc requires a target with CTA register reconfiguration, such as sm_90a

The generated CUDA code is compiled using an NVCC command containing:

-arch=sm_110a

The failure occurs across essentially all tested GEMM autotuning configurations, independent of the particular GEMM tile sizes.

For example, configurations such as:

block_M=64
block_N=64
block_K=32
num_stages=2 or 3
threads=128 or 256

fail at the same CTA register-reconfiguration check.

Workaround / Isolation

TileLang provides an option to disable warp-group register redistribution while keeping the warp-specialized producer-consumer pipeline enabled:

T.disable_warp_group_reg_alloc()

After enabling this option, the previous compilation failures disappear and the generated kernels successfully proceed past compilation.

This appears to isolate the issue specifically to the setmaxnreg / CTA register-reconfiguration path rather than to GEMM layout inference, software pipelining, TMA, or the autotuning configurations themselves.

In other words, the behavior is approximately:

Warp specialization + register redistribution
    -> compilation failure on sm_110a

Warp specialization + register redistribution disabled
    -> compilation succeeds

Relevant CUTLASS Code

The current include/cutlass/arch/reg_reconfig.h contains feature detection similar to:

#ifndef CUDA_CTA_RECONFIG_ACTIVATED
  #if defined(__CUDA_ARCH__) && __CUDACC_VER_MAJOR__ >= 12 && (
         (__CUDA_ARCH__ ==  900 && defined(__CUDA_ARCH_FEAT_SM90_ALL))
      || (__CUDA_ARCH__ == 1000 && defined(__CUDA_ARCH_FEAT_SM100_ALL))
      || (__CUDA_ARCH__ == 1010 && defined(__CUDA_ARCH_FEAT_SM101_ALL))
      || (__CUDA_ARCH__ == 1030 && defined(__CUDA_ARCH_FEAT_SM103_ALL))
      || (__CUDA_ARCH__ == 1200 && defined(__CUDA_ARCH_FEAT_SM120_ALL))
      || (__CUDA_ARCH__ == 1210 && defined(__CUDA_ARCH_FEAT_SM121_ALL))
    )
    #define CUDA_CTA_RECONFIG_ACTIVATED 1
  #endif

  ...

  #if defined(__CUDA_ARCH__) && __CUDACC_VER_MAJOR__ >= 12 && (
         (__CUDA_ARCH__ == 1070 && CUDA_ARCH_FAMILY(1070))
      || (__CUDA_ARCH__ == 1070 && defined(__CUDA_ARCH_FEAT_SM107_ALL))
    )
    #define CUDA_CTA_RECONFIG_ACTIVATED 1
  #endif
#endif

There does not appear to be a corresponding SM110/SM110a case.

Expected Behavior / Question

Is the omission of SM110a from CUDA_CTA_RECONFIG_ACTIVATED intentional?

If CTA register reconfiguration and the following PTX instructions are supported on the SM110a architecture-specific target:

setmaxnreg.inc.sync.aligned.u32
setmaxnreg.dec.sync.aligned.u32

would it be appropriate to enable CUDA_CTA_RECONFIG_ACTIVATED for SM110a in reg_reconfig.h as well?

Alternatively, if SM110a requires a different architecture feature macro or compilation mode, could you clarify the intended way to detect CTA register-reconfiguration support on Jetson AGX Thor?

For example, should the check depend on an SM110-specific architecture feature macro, or does SM110a require a particular compute_110a / sm_110a code-generation configuration?

Performance Impact

Disabling warp-group register redistribution is a functional workaround, but it is not necessarily performance-neutral.

The warp-specialized kernel still separates producer and consumer warp groups, but it can no longer perform register redistribution such as:

Producer warp group
    -> setmaxnreg.dec
    -> releases register budget

Consumer warp group
    -> setmaxnreg.inc
    -> receives additional register budget

This may increase register pressure or spilling in the GEMM consumer and therefore potentially reduce performance.

For this reason, supporting the intended CTA register-reconfiguration path on SM110a would be preferable to permanently disabling this optimization.

Additional Information

I can provide:

• the complete NVCC command line,
• nvcc --version,
• the generated CUDA source,
• TileLang version / commit,
• and a minimal reproducer,

if useful for reproducing the issue.

Start by reading include/cutlass/arch/reg_reconfig.h and reproduce the reported failure with the sm_110a NVCC target, using the described TileLang warp-specialized GEMM path. Check the available SM110a architecture feature macros and CTA register-reconfiguration support; done means the intended detection is established and the relevant compilation path is covered without the workaround.