[Build] update version of "cutlass"

[mc-nv]

Describe the issue

Triton team facing build issue trying to compile ONNX Runtime.
That issue observed against CUDA 12.4 and "cutlass" v3.1.0.

Note

1. Were able to suppress that warning with --compile_no_warning_as_error
2. Haven't seen issue trying compile against latest version of "cutlass"

Urgency

No response

Target platform

WIN32

Build script

build.bat --cmake_generator "Visual Studio 17 2022" --config Release --cmake_extra_defines "CMAKE_CUDA_ARCHITECTURES=60;61;70;75;80;86;90" --skip_submodule_sync --parallel --build_shared_lib --update --build --build_dir /workspace/build --use_cuda --cuda_version "12.4" --cuda_home "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4" --cudnn_home "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4" --use_tensorrt --tensorrt_home "/tensorrt"

Error / output

         C:\workspace\build\Release\_deps\cutlass-src\include\cute/arch/cluster_sm90.hpp(101): error #940-D: missing return statement at end of non-void function "cute::cluster_grid_dims" [C:\workspace\build\Release\onnxruntime_providers_cuda.vcxproj] [C:\tmp\tritonbuild\onnxruntime\build\ort_target.vcxproj]
         C:\workspace\build\Release\_deps\cutlass-src\include\cute/arch/cluster_sm90.hpp(120): error #940-D: missing return statement at end of non-void function "cute::cluster_id_in_grid" [C:\workspace\build\Release\onnxruntime_providers_cuda.vcxproj] [C:\tmp\tritonbuild\onnxruntime\build\ort_target.vcxproj]
        

Visual Studio Version

BUILDTOOLS_VERSION:17.9.34622.214 CMAKE_VERSION:3.27.1 CUDA_VERSION:12.4.0 CUDNN_VERSION:9.0.0.312 PYTHON_VERSION:3.8.10 TENSORRT_VERSION:8.6.1.6 VCPGK_VERSION:2023.11.20

GCC / Compiler Version

No response

[jywu-msft]

+@tianleiwu @yufenglee

[jywu-msft]

yes, I encountered this as well. there seems to be a bug in cutlass 3.1
see: https://github.com/NVIDIA/cutlass/blob/6f47420213f757831fae65c686aa471749fa8d60/include/cute/arch/cluster_sm90.hpp#L97

latest cutlass seems to have it fixed
https://github.com/NVIDIA/cutlass/blob/ffa34e70756b0bc744e1dfcc115b5a991a68f132/include/cute/arch/cluster_sm90.hpp#L118

however when I update to cutlass 3.4.1 and retry build, I encounter other build errors
e.g.
C:\ort-118-cuda124-trt10\Release_deps\cutlass-src\include\cutlass/epilogue/warp/tile_iterator_tensor
_op_mixed.h(109): error : expression must have a constant value [C:\ort-118-cuda124-trt10\Release\onn
xruntime_providers_cuda.vcxproj]

which version of cutlass did you use to get a successful build?

[mc-nv]

"cutlass" compilation was against main branch.

[jywu-msft]

just to confirm, you had to update cutlass to main AND also suppress warnings as error to get the build to succeed?

[github-actions]

This issue has been automatically marked as stale due to inactivity and will be closed in 30 days if no further activity occurs. If further support is needed, please provide an update and/or more details.

[snnn]

Do we need to upgrade cutlass version?

[tianleiwu]

Yes, it need upgrade cutlass version to 3.5 and need some code change in ORT. I'm working on that.