search

NVIDIA / cutlass

CUDA Templates for Linear Algebra Subroutines
Other
5.64k stars 961 forks source link

[BUG] Unable to build against CUDA 12.4 without #1403

Closed mc-nv closed 3 months ago

mc-nv commented 8 months ago

Describe the bug Unable to compile cutlass source code against CUDA 12.4

Steps/Code to reproduce bug

PS C:\workspace\cutlass> history

  Id CommandLine
  -- -----------
   1 md workspace
   2 cd workspace
   3 git clone v3.1.0 https://github.com/NVIDIA/cutlass.git
   4 cd cutlass
   5 cmake -B build
   6 cmake --build build

Getting error: 

error #940-D: missing return statement at end of non-void function "cute::cluster_grid_dims" 
error #940-D: missing return statement at end of non-void function "cute::cluster_id_in_grid"

Expected behavior Shouldn't compile without issues.

Environment details (please complete the following information): Docker, Bare metal 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

Additional context https://github.com/microsoft/onnxruntime/issues/19891

thakkarV commented 8 months ago

Hi! This is for CUTLASS version 3.1 which was released quite a few months ago (before the release of CUDA 12.4). Are you able to repro this with CUTLASS 3.4? It seems like you are also building on Windows, and our support for windows builds has improved quite a bit since 3.1.

mc-nv commented 8 months ago

My latest build on Windows within same environment stuck for several hours against main. And trow errors like: 

C:\workspace\cutlass\include\cute/layout.hpp(746): catastrophic error : out of memory [C:\workspace\cutlass\b
uild\tools\library\cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.vcxproj]
        return bw_coalesce<I-1>(old_shape, old_stride, new_shape, new_stride);
        ^

  1 catastrophic error detected in the compilation of "C:/workspace/cutlass/build/tools/library/cutlass_libra
  ry_gemm_sm90_void_s64x128x16gemm_f16_objs.unity.daa0ddf64e7b.cu".
  Compilation terminated.
  cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.unity.daa0ddf64e7b.cu
C:\BuildTools\MSBuild\Microsoft\VC\v170\BuildCustomizations\CUDA 12.4.targets(799,9): error MSB3721: The comm
and ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\bin\nvcc.exe"  --use-local-env -ccbin "C:\Buil
dTools\VC\Tools\MSVC\14.39.33519\bin\HostX64\x64" -x cu   -IC:\workspace\cutlass\include -IC:\workspace\cutla
ss\build\include -I\include -I\examples -IC:\workspace\cutlass\tools\library\include -IC:\workspace\cutlass\t
ools\util\include -IC:\workspace\cutlass\tools\library\src -IC:\workspace\cutlass\build\tools\library\include
 -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\include" -I"C:\Program Files\NVIDIA GPU Computin
g Toolkit\CUDA\v12.4\include"     --keep-dir cutlass_.55939AC6\x64\Debug  -maxrregcount=0   --machine 64 --co
mpile -cudart static -std=c++17 --generate-code=arch=compute_70,code=[sm_70] --generate-code=arch=compute_70,
code=[compute_70] --generate-code=arch=compute_72,code=[sm_72] --generate-code=arch=compute_72,code=[compute_
72] --generate-code=arch=compute_75,code=[sm_75] --generate-code=arch=compute_75,code=[compute_75] --generate
-code=arch=compute_80,code=[sm_80] --generate-code=arch=compute_80,code=[compute_80] --generate-code=arch=com
pute_86,code=[sm_86] --generate-code=arch=compute_86,code=[compute_86] --generate-code=arch=compute_87,code=[
sm_87] --generate-code=arch=compute_87,code=[compute_87] --generate-code=arch=compute_89,code=[sm_89] --gener
ate-code=arch=compute_89,code=[compute_89] --generate-code=arch=compute_90,code=[sm_90] --generate-code=arch=
compute_90,code=[compute_90] --generate-code=arch=compute_90a,code=[sm_90a] --generate-code=arch=compute_90a,
code=[compute_90a] --expt-relaxed-constexpr -Xcompiler="/EHsc /Zc:__cplusplus /bigobj -Zi -Ob0 /wd4819 /fp:st
rict" -g  -D_WINDOWS -DCUTLASS_VERSIONS_GENERATED -DCUTLASS_ENABLE_TENSOR_CORE_MMA=1 -DCUTLASS_TEST_LEVEL=0 -
DCUTLASS_TEST_ENABLE_CACHED_RESULTS=1 -DCUTLASS_CONV_UNIT_TEST_RIGOROUS_SIZE_ENABLED=1 -DCUTLASS_DEBUG_TRACE_
LEVEL=0 -D"CMAKE_INTDIR=\"Debug\"" -D_MBCS -D"CMAKE_INTDIR=\"Debug\"" -Xcompiler "/EHsc /W3 /nologo /Od /FS /
Zi /RTC1 /MDd /GR" -Xcompiler "/Fdcutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.dir\Debug\cutlass_li
brary_gemm_sm90_void_s64x128x16gemm_f16_objs.pdb" -o cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.d
ir\Debug\cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.unity.daa0ddf64e7b.obj "C:\workspace\cutlass\
build\tools\library\cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.unity.daa0ddf64e7b.cu"" exited wit
h code 1. [C:\workspace\cutlass\build\tools\library\cutlass_library_gemm_sm90_void_s64x128x16gemm_f16_objs.vc
xproj]
thakkarV commented 8 months ago

That's an out of memory error so likely an issue with the compiler or the system used to build the kernels? CC @mhoemmen

mc-nv commented 8 months ago

I got following error output from ONNX Runtime build trying to engage the "cutlass" as a CMake submodule against CUDA 12.4: 

C:\workspace\Release\_deps\cutlass-src\include\cutlass/gemm/threadblock/default_mma_core_sm80.h(2495): error : expression mus
t have a constant value [C:\ort-118-cuda124-trt10-latestCutlass\Release\onnxruntime_providers_cuda.vcxproj]
      static const int LaneN = cutlass::const_min(numElementsB, ThreadTileN);
                               ^
  C:\workspace\Release\_deps\cutlass-src\include\cutlass/gemm/threadblock/default_mma_core_sm80.h(2495): note
thakkarV commented 7 months ago

CC @hwu36 and @mhoemmen

mhoemmen commented 7 months ago

@mc-nv Thanks for the error report! We'll take a look at this.

yf711 commented 7 months ago

Hi @mhoemmen any update on this?

mhoemmen commented 7 months ago

@yf711 wrote:

I got following error output from ONNX Runtime build trying to engage the "cutlass" as a CMake submodule against CUDA 12.4:

I'm not sure what it means to "engage the 'cutlass' as a CMake submodule against CUDA 12.4."

Could you please post all the CMake options that were given to CUTLASS? Without the list of CMake options, it will be a lot harder for us to try to reproduce this.

github-actions[bot] commented 6 months ago

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

chilo-ms commented 5 months ago

Could you please post all the CMake options that were given to CUTLASS? Without the list of CMake options, it will be a lot harder for us to try to reproduce this.

Hi @mhoemmen

OnnxRuntime's CUDA EP includes cutlass header directories to its cmake target:

    ...
    include(cutlass)
    target_include_directories(${target} PRIVATE ${cutlass_SOURCE_DIR}/include ${cutlass_SOURCE_DIR}/examples ${cutlass_SOURCE_DIR}/tools/util/include)
    ...

https://github.com/microsoft/onnxruntime/blob/main/cmake/external/cutlass.cmake https://github.com/microsoft/onnxruntime/blob/main/cmake/onnxruntime_providers_cuda.cmake#L214

So some of its header files can include cutlass header files moe_gemm_kernels_template.h

The compiler errors were shown when building the CUDA EP target. You can either see the compiler flags from the onnxruntime_providers_cuda.cmake above or command line log below.

With cutlass 3.1.0:

C:/Users/lochi/repos/onnxruntime/build/Windows/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:\Users\lochi\repos\onnxruntime\build\Windows\Release\onnxruntime_prov
iders_cuda.vcxproj]
    }
    ^

  Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

C:/Users/lochi/repos/onnxruntime/build/Windows/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:\Users\lochi\repos\onnxruntime\build\Windows\Release\onnxruntime_pro
viders_cuda.vcxproj]
    }
    ^

  2 errors detected in the compilation of "C:/Users/lochi/repos/onnxruntime/onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_gemm_kernels_fp32_fp32.cu".
  moe_gemm_kernels_fp32_fp32.cu
C:\Program Files\Microsoft Visual Studio\2022\Enterprise\MSBuild\Microsoft\VC\v170\BuildCustomizations\CUDA 12.4.targets(799,9): error MSB3721: The command ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\bin\nvcc.exe"  --use-local-env -ccbin "C:\Program Files\Microsof
t Visual Studio\2022\Enterprise\VC\Tools\MSVC\14.40.33807\bin\HostX64\x64" -x cu   -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\utf8_range-src" -IC:\Users\lochi\repos\onnxruntime\include\onnxruntime -IC:\Users\lochi\repos\onnxruntime\include\onnxruntime\core\s
ession -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\pytorch_cpuinfo-src\include" -IC:\Users\lochi\repos\onnxruntime\build\Windows\Release -IC:\Users\lochi\repos\onnxruntime\onnxruntime -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\abseil_cpp-
src" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\safeint-src" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\gsl-src\include" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\date-src\include" -I"C:\Users\lochi\repos\onnxruntim
e\build\Windows\Release\_deps\onnx-src" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\onnx-build" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\protobuf-src\src" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\flatbuffers-src\i
nclude" -I"C:\Users\lochi\Downloads\cudnn-windows-x86_64-8.9.7.29_cuda12-archive\cudnn-windows-x86_64-8.9.7.29_cuda12-archive\include" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\cutlass-src\include" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_
deps\cutlass-src\examples" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\cutlass-src\tools\util\include" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\eigen-src" -I"C:\Users\lochi\tensorrt\TensorRT-10.0.1.6.Windows10.win10.cuda-12.4\TensorRT-1
0.0.1.6\include" -I"C:\Users\lochi\repos\onnxruntime\build\Windows\Release\_deps\mp11-src\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\include"     --keep-dir onnxrunt.B5EE5B2F\x64\Relea
se  -maxrregcount=0   --machine 64 --compile -cudart shared --expt-relaxed-constexpr --Werror default-stream-launch -Xcudafe --diag_suppress=bad_friend_decl -Xcudafe --diag_suppress=unsigned_compare_with_zero -Xcudafe --diag_suppress=expr_has_no_effect -include algorithm -std=c+
+17 --generate-code=arch=compute_60,code=[compute_60,sm_60] --generate-code=arch=compute_61,code=[compute_61,sm_61] --generate-code=arch=compute_70,code=[compute_70,sm_70] --generate-code=arch=compute_75,code=[compute_75,sm_75] --generate-code=arch=compute_80,code=[compute_80,sm
_80] --generate-code=arch=compute_86,code=[compute_86,sm_86] --generate-code=arch=compute_90,code=[compute_90,sm_90] -Xcudafe --diag_suppress=conversion_function_not_usable --threads 1 -Werror all-warnings -Xcompiler="/EHsc -Ob2 /utf-8 /sdl /experimental:external /external:W0 /e
xternal:templates- /external:IC:/Users/lochi/repos/onnxruntime/cmake /external:IC:/Users/lochi/repos/onnxruntime/build/Windows/Release /wd4251 /wd4201 /wd4324 /wd5054 /w15038 /wd4251 /wd4201 /wd4324 /wd5054 /w15038 /wd4834 /wd4127"   -D_WINDOWS -DNDEBUG -DCPUINFO_SUPPORTED_PLATF
ORM=1 -DEIGEN_USE_THREADS -DDISABLE_CUSPARSE_DEPRECATED -DPLATFORM_WINDOWS -DNOGDI -DNOMINMAX -D_USE_MATH_DEFINES -D_SILENCE_ALL_CXX17_DEPRECATION_WARNINGS -DUSE_CUDA=1 -DUSE_MEMORY_EFFICIENT_ATTENTION=1 -DUSE_TENSORRT=1 -DONLY_C_LOCALE=0 -DONNX_NAMESPACE=onnx -DONNX_ML=1 -DONNX
_USE_LITE_PROTO=1 -D__ONNX_NO_DOC_STRINGS -DWIN32_LEAN_AND_MEAN -DORT_ENABLE_STREAM -DEIGEN_MPL2_ONLY -DEIGEN_HAS_CONSTEXPR -DEIGEN_HAS_VARIADIC_TEMPLATES -DEIGEN_HAS_CXX11_MATH -DEIGEN_HAS_CXX11_ATOMIC -DEIGEN_STRONG_INLINE=inline -D_SILENCE_EXPERIMENTAL_FILESYSTEM_DEPRECATION_
WARNING=1 -D"CMAKE_INTDIR=\"Release\"" -Donnxruntime_providers_cuda_EXPORTS -D_WINDLL -D_MBCS -DEIGEN_HAS_C99_MATH -DCPUINFO_SUPPORTED -DNDEBUG -DCPUINFO_SUPPORTED_PLATFORM=1 -DEIGEN_USE_THREADS -DDISABLE_CUSPARSE_DEPRECATED -DPLATFORM_WINDOWS -DNOGDI -DNOMINMAX -D_USE_MATH_DEFI
NES -D_SILENCE_ALL_CXX17_DEPRECATION_WARNINGS -DUSE_CUDA=1 -DUSE_MEMORY_EFFICIENT_ATTENTION=1 -DUSE_TENSORRT=1 -DONLY_C_LOCALE=0 -DONNX_NAMESPACE=onnx -DONNX_ML=1 -DONNX_USE_LITE_PROTO=1 -D__ONNX_NO_DOC_STRINGS -DWIN32_LEAN_AND_MEAN -DORT_ENABLE_STREAM -DEIGEN_MPL2_ONLY -DEIGEN_
HAS_CONSTEXPR -DEIGEN_HAS_VARIADIC_TEMPLATES -DEIGEN_HAS_CXX11_MATH -DEIGEN_HAS_CXX11_ATOMIC -DEIGEN_STRONG_INLINE=inline -D_SILENCE_EXPERIMENTAL_FILESYSTEM_DEPRECATION_WARNING=1 -D"CMAKE_INTDIR=\"Release\"" -Donnxruntime_providers_cuda_EXPORTS -Xcompiler "/EHsc /W4 /nologo /O2
/FS   /MD /GR" -Xcompiler "/Fdonnxruntime_providers_cuda.dir\Release\vc143.pdb" -o onnxruntime_providers_cuda.dir\Release\moe_gemm_kernels_fp32_fp32.obj "C:\Users\lochi\repos\onnxruntime\onnxruntime\contrib_ops\cuda\moe\ft_moe\moe_gemm_kernels_fp32_fp32.cu"" exited with code 2.
[C:\Users\lochi\repos\onnxruntime\build\Windows\Release\onnxruntime_providers_cuda.vcxproj]

With cutlass 3.5.0:

mhoemmen commented 5 months ago

A colleague has been communicating offline to see if setting the /Zc:__cplusplus flag helps. I'm awaiting news on that.

tianleiwu commented 5 months ago

/Zc:__cplusplus flag resolved the issue. Thanks.

mhoemmen commented 5 months ago

Thanks @tianleiwu ! : - )

@hwu36 It looks like we can close this issue.

github-actions[bot] commented 4 months ago

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.