[BUG] CUTLASS_ARCH_MMA_SM100A_ENABLED in ./examples/82_blackwell_distributed_gemm/82_blackwell_distributed_gemm.cu

[slove2857]

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug
It outputs 'This example must be compiled with sm100a and CUDA Toolkit 12.8 or later' even after following the cmake command 'cmake $PATH -DCUTLASS_NVCC_ARCHS="100a" -DCUTLASS_ENABLE_GDC_FOR_SM100=1
'.

Manually output the following ENVs
--- Environment Check Status ---
CUTLASS_ARCH_MMA_SM100A_ENABLED: NOT DEFINED
CUDACC_VER_MAJOR: 13
CUDACC_VER_MINOR: 0
Detected CUDA Toolkit Version: 13.0

I manually comment all CUTLASS_ARCH_MMA_SM100A_ENABLED in the cuda code, and then it can work.

So I guess CUTLASS_ARCH_MMA_SM100A_ENABLED is unnecessary if we have cuda >= 12.8? If this is true, any chance to remove this flag so that it does not block running from raw repo?

[AIInfra-lvz]

Hi! The key might be the SM version, while the CUDA version is just a requirement. I noticed that instructions such as tcgen05 are enabled or disabled by the macro CUTLASS_ARCH_MMA_SM100A_ENABLED. You can find more details in the files cute/arch/config.h and cute/arch/config.hpp. Hope this helps!

[alihassanijr]

Thank you for bringing this to our attention.

Indeed CUTLASS_ARCH_MMA_SM100A_ENABLED is wrong, but mainly because it is a macro that is only effective on device code, not host.

Like other Blackwell examples (and Hopper too for that matter), the kernel instantiations must be guarded with CUTLASS_ARCH_MMA_SM100_SUPPORTED (not CUTLASS_ARCH_MMA_SM100A_ENABLED) to make sure the compiler is enabling their respective TC instructions.

The host code needs to be separately guarded to make sure the (relatively) new CUDA graph APIs used are available.

#2713 should address this issue.