[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_90aThe generated CUDA code is compiled using an NVCC command containing:
-arch=sm_110aThe 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 256fail 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 succeedsRelevant 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
#endifThere 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.u32would 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 budgetThis 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.
Source: NVIDIA/cutlass