Patches to CK makefile

[zjing14]

Refer to: https://github.com/microsoft/onnxruntime/blob/main/cmake/patches/composable_kernel/Fix_Clang_Build.patch Need the following changes in CK cmakefile

Line: 9

 # Check support for CUDA/HIP in Cmake
-project(composable_kernel)
+project(composable_kernel LANGUAGES CXX HIP)

Line: 46 -link_libraries(hip::device)

Line: 80

+    set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
     add_library(${INSTANCE_NAME} OBJECT ${ARGN})
     target_compile_features(${INSTANCE_NAME} PUBLIC)
+    target_compile_definitions(${INSTANCE_NAME} PRIVATE "__HIP_PLATFORM_AMD__=1" "__HIP_PLATFORM_HCC__=1")

[cloudhan]

Ping me if you have any question.

This following is also some very important things you should take care of. https://github.com/microsoft/onnxruntime/blob/7b31bcda2e9cd45c709e3fe31a544297db37ea3c/cmake/CMakeLists.txt#L255-L262

-mllvm -amdgpu-early-inline-all=true is propagated from some underlying targets, you should make sure they are properly added. Otherwise, there will be huge perf cliff for a subset of ck kernels.

[zjing14]

@illsilin Could you take care of it?

[illsilin]

As soon as I replace "project(composable_kernel)" with "project(composable_kernel LANGUAGES CXX HIP)", I get the following errors with our standard cmake command:

root@rocm-framework-11:/composable_kernel/build# cmake -DBUILD_DEV=OFF -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -DGPU_TARGETS="gfx908;gfx90a" -DCMAKE_PREFIX_PATH=/opt/rocm .. CMake Error: Could not find cmake module file: CMakeDetermineHIPCompiler.cmake CMake Error: Error required internal CMake variable not set, cmake may not be built correctly. Missing variable is: CMAKE_HIP_COMPILER_ENV_VAR CMake Error: Error required internal CMake variable not set, cmake may not be built correctly. Missing variable is: CMAKE_HIP_COMPILER CMake Error: Could not find cmake module file: /composable_kernel/build/CMakeFiles/3.16.3/CMakeHIPCompiler.cmake CMake Error at CMakeLists.txt:4 (project): No CMAKE_HIP_COMPILER could be found.

We can add "-DCMAKE_HIP_COMPILER=/opt/rocm/bin/hipcc" to the cmake options, but what is the value for the "CMAKE_HIP_COMPILER_ENV_VAR"?

[jithunnair-amd]

As soon as I replace "project(composable_kernel)" with "project(composable_kernel LANGUAGES CXX HIP)", I get the following errors with our standard cmake command:

root@rocm-framework-11:/composable_kernel/build# cmake -DBUILD_DEV=OFF -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -DGPU_TARGETS="gfx908;gfx90a" -DCMAKE_PREFIX_PATH=/opt/rocm .. CMake Error: Could not find cmake module file: CMakeDetermineHIPCompiler.cmake CMake Error: Error required internal CMake variable not set, cmake may not be built correctly. Missing variable is: CMAKE_HIP_COMPILER_ENV_VAR CMake Error: Error required internal CMake variable not set, cmake may not be built correctly. Missing variable is: CMAKE_HIP_COMPILER CMake Error: Could not find cmake module file: /composable_kernel/build/CMakeFiles/3.16.3/CMakeHIPCompiler.cmake CMake Error at CMakeLists.txt:4 (project): No CMAKE_HIP_COMPILER could be found.

We can add "-DCMAKE_HIP_COMPILER=/opt/rocm/bin/hipcc" to the cmake options, but what is the value for the "CMAKE_HIP_COMPILER_ENV_VAR"?

@illsilin I ran into this in one of my MIOpen builds for the rocm-rel-6.2 branch of MIOpen, in an environment with CMake 3.18.4 (newer than in your snippet). How was this fixed? If CK now requires a newer CMake version to find CMakeDetermineHIPCompiler.cmake, should the minimum CMake version be bumped up here: https://github.com/ROCm/composable_kernel/blob/8c90f25be3fc3dcf5ef2475107bd7130b7c9c7ae/CMakeLists.txt#L1?

[illsilin]

OK, I just checked our latest main cmake file. Looks like we already have all the compiler options mentioned above except the "-mllvm -amdgpu-early-inline-all=true". I can try to add it and see what the performance impact is going to be. If there are no degradations, i'll merge it.

[illsilin]

Actually, it looks like there are some issues with that flag in the latest compiler versions: https://github.com/llvm/llvm-project/issues/86332

[illsilin]

OK, when i build with verbose flag on, i can actually see that flag present. It's apparently one of the HIP default flags. So closing this issue.

[jithunnair-amd]

@illsilin Can you please let me know which CMake version you successfully built CK with?