Enable tcgen05 blockscaled ops on Thor SM110

Edge-LLM NvFP4 MoE CuTeDSL kernels on Thor use tcgen05 blockscaled MMA and SMEM-to-TMEM scale-factor copies. The existing checks only admitted the SM100/SM103 paths, so source-built CuTeDSL rejected SM110.

Admit Thor's blockscaled MMA arch aliases sm_101a and sm_110a, and allow the SM110f family for S2T tcgen05 copy ops.

Validation:

- git diff --check

- python3 -m py_compile python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py

- DKG grouped_blockscaled_gemm.py documented 4-group example on Thor SM110: PASS

- Edge-LLM nvfp4_moe AOT for sm_110/aarch64: 12/12 variants PASS

Author: xiangg-nv

python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py

@@ -786,16 +786,20 @@ class _S2TCopyBase(CopyOp):
 
     :param cta_group: Cooperative Thread Array (CTA) group configuration
     :type cta_group: CtaGroup
-    :raises OpError: If the current architecture is not SM100f family or if invalid parameters are provided
+    :raises OpError: If the current architecture is not SM100f or SM110f
+        family or if invalid parameters are provided
     """
 
     cta_group: CtaGroup
 
     def __post_init__(self) -> None:
         # Arch verification
         arch = BaseDSL._get_dsl().get_arch_enum()
-        if not arch.is_family_of(Arch.sm_100f):
-            supported = Arch.filter(lambda a: a.is_family_of(Arch.sm_100f))
+        # S2T tcgen05 copy encodings are valid on both SM100 and Thor SM110.
+        if not (arch.is_family_of(Arch.sm_100f) or arch.is_family_of(Arch.sm_110f)):
+            supported = Arch.filter(
+                lambda a: a.is_family_of(Arch.sm_100f) or a.is_family_of(Arch.sm_110f)
+            )
             raise OpError(
                 self,
                 f"expects arch to be one of {supported}, but got {arch}",

python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py

@@ -386,7 +386,9 @@ class BlockScaledMmaOp(Tcgen05MmaOp):
 
     admissible_archs = [
         Arch.sm_100a,
+        Arch.sm_101a,
         Arch.sm_103a,
+        Arch.sm_110a,
     ]
 
     def __post_init__(self) -> None: