[QST] Can't compile 57_hopper_grouped_gemm example

[andrewkchan]

Hi, I'm getting some compile errors when trying to build 57_hopper_grouped_gemm on my machine. Wondering if there's some include path that's missing for some reason in my setup? Here are the errors I get after doing cmake in the project root then running make in the example folder:

$ ~/achan/cutlass/examples/57_hopper_grouped_gemm$ make
Building CUDA object examples/57_hopper_grouped_gemm/CMakeFiles/57_hopper_grouped_gemm.dir/57_hopper_grouped_gemm.cu.o
/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(223): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, float_e2m3_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                       ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(224): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, float_e3m2_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                       ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(225): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B" is undefined
    if constexpr (is_same_v<T, float_e2m1_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B;} else
                                                       ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(226): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B" is undefined
    if constexpr (is_same_v<T, cutlass::detail::float_e2m1_unpacksmem_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B;} else
                                                                                   ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(227): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, cutlass::detail::float_e2m3_unpacksmem_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                                                   ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(228): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, cutlass::detail::float_e3m2_unpacksmem_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                                                   ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(229): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, detail::type_erased_dynamic_float6_unpacksmem_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                                                          ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(230): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B" is undefined
    if constexpr (is_same_v<T, type_erased_dynamic_float6_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B;} else
                                                                       ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(231): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B" is undefined
    if constexpr (is_same_v<T, detail::type_erased_dynamic_float4_unpacksmem_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B;} else
                                                                                          ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(232): error: identifier "CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B" is undefined
    if constexpr (is_same_v<T, type_erased_dynamic_float4_t>) { return CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B; } else
                                                                       ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(260): error: identifier "CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B" is undefined
          case SmemSwizzleBase::SWIZZLE_BASE_32B: return CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B;
                                                         ^

/mnt/large_shared/testuser/achan/cutlass/include/cute/arch/copy_sm90_desc.hpp(261): error: identifier "CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B" is undefined
          case SmemSwizzleBase::SWIZZLE_BASE_64B: return CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B;
                                                         ^

12 errors detected in the compilation of "/mnt/large_shared/testuser/achan/cutlass/examples/57_hopper_grouped_gemm/57_hopper_grouped_gemm.cu".
make[2]: *** [examples/57_hopper_grouped_gemm/CMakeFiles/57_hopper_grouped_gemm.dir/build.make:76: examples/57_hopper_grouped_gemm/CMakeFiles/57_hopper_grouped_gemm.dir/57_hopper_grouped_gemm.cu.o] Error 2
make[1]: *** [CMakeFiles/Makefile2:47326: examples/57_hopper_grouped_gemm/CMakeFiles/57_hopper_grouped_gemm.dir/all] Error 2
make: *** [Makefile:166: all] Error 2

Here's the output of nvcc --version:

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Tue_Oct_29_23:50:19_PDT_2024
Cuda compilation tools, release 12.6, V12.6.85
Build cuda_12.6.r12.6/compiler.35059454_0

[zhyncs]

same issue

[LANSHANGH]

same issue

[ANIKET-SHIVAM]

You can use v3.7.0 release. Or use CUDA 12.8 for now. Will be fixed to work with older CUDA versions very soon, it's not specific to just this example.