Merge pull request #28 from ROCm/ck_update

CK update for staging build as well as a general pull for improvements

cgmanifests/generated/cgmanifest.json

@@ -336,7 +336,7 @@
       "component": {
         "type": "git",
         "git": {
-          "commitHash": "5356c4a943a35e74d7cdc69486afcb8703b9a59a",
+          "commitHash": "2fd6c6d4e7d7c458bd44f4df49271b95ce745b20",
           "repositoryUrl": "https://github.com/ROCmSoftwarePlatform/composable_kernel.git"
         },
         "comments": "composable_kernel"

cmake/deps.txt

@@ -56,4 +56,4 @@ tensorboard;https://github.com/tensorflow/tensorboard/archive/373eb09e4c5d2b3cc2
 cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v3.1.0.zip;757f90a795034a89d4f48a79d1f009f7a04c8dee
 utf8_range;https://github.com/protocolbuffers/utf8_range/archive/72c943dea2b9240cd09efde15191e144bc7c7d38.zip;9925739c9debc0efa2adcb194d371a35b6a03156
 extensions;https://github.com/microsoft/onnxruntime-extensions/archive/94142d8391c9791ec71c38336436319a2d4ac7a0.zip;4365ac5140338b4cb75a39944a4be276e3829b3c
-composable_kernel;https://github.com/ROCmSoftwarePlatform/composable_kernel/archive/5356c4a943a35e74d7cdc69486afcb8703b9a59a.zip;522382c2af437e09124287e5879ab64af5b2e299
+composable_kernel;https://github.com/ROCmSoftwarePlatform/composable_kernel/archive/2fd6c6d4e7d7c458bd44f4df49271b95ce745b20.zip;37b256750556bd19c08d81703f27fea3fea4fe33

cmake/patches/composable_kernel/Fix_Clang_Build.patch

@@ -1,17 +1,17 @@
 diff --git a/CMakeLists.txt b/CMakeLists.txt
-index 04674124c..12e8b8b00 100644
+index bdeba33e..afe406b3 100644
 --- a/CMakeLists.txt
 +++ b/CMakeLists.txt
-@@ -19,7 +19,7 @@ endif()
+@@ -23,7 +23,7 @@ endif()
 
  set(version 1.1.0)
  # Check support for CUDA/HIP in Cmake
--project(composable_kernel VERSION ${version})
+-project(composable_kernel VERSION ${version} LANGUAGES CXX)
 +project(composable_kernel VERSION ${version} LANGUAGES CXX HIP)
+ include(CTest)
 
  list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
-
-@@ -173,27 +173,6 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON)
+@@ -206,27 +206,6 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON)
  set(CMAKE_CXX_EXTENSIONS OFF)
  message("CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")
 
@@ -39,27 +39,20 @@ index 04674124c..12e8b8b00 100644
  ## HIP
  find_package(HIP REQUIRED)
  # Override HIP version in config.h, if necessary.
-@@ -215,8 +194,6 @@ if( DEFINED CK_OVERRIDE_HIP_VERSION_PATCH )
+@@ -248,12 +227,6 @@ if( DEFINED CK_OVERRIDE_HIP_VERSION_PATCH )
      message(STATUS "CK_HIP_VERSION_PATCH overriden with ${CK_OVERRIDE_HIP_VERSION_PATCH}")
  endif()
  message(STATUS "Build with HIP ${HIP_VERSION}")
 -link_libraries(hip::device)
--add_compile_definitions(__HIP_PLATFORM_HCC__=1)
+-if(CK_hip_VERSION VERSION_GREATER_EQUAL 6.0.23494)
+-    add_compile_definitions(__HIP_PLATFORM_AMD__=1)
+-else()
+-    add_compile_definitions(__HIP_PLATFORM_HCC__=1)
+-endif()
 
  ## tidy
  include(EnableCompilerWarnings)
-@@ -376,7 +353,9 @@ if(BUILD_DEV)
-     add_compile_options(-Werror -Weverything)
- endif()
- #add flags to reduce the size of binaries
--add_compile_options(-Oz -flto=thin)
-+# -flto requires ORT to use a linker that support LTO and -flto flag shoud be passed to linker together.
-+# add_compile_options(-Oz -flto=thin)
-+add_compile_options(-Oz)
- message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}")
-
- add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR})
-@@ -482,11 +461,3 @@ rocm_install(FILES
+@@ -520,11 +493,3 @@ rocm_install(FILES
 
  set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE")
  set(CPACK_RPM_PACKAGE_LICENSE "MIT")
@@ -72,10 +65,10 @@ index 04674124c..12e8b8b00 100644
 -    HEADER_ONLY
 -)
 diff --git a/library/src/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/CMakeLists.txt
-index 9cb5d0e9a..141a46f3d 100644
+index 0a12e1c4..dcdaaf5a 100644
 --- a/library/src/tensor_operation_instance/gpu/CMakeLists.txt
 +++ b/library/src/tensor_operation_instance/gpu/CMakeLists.txt
-@@ -44,8 +44,14 @@ function(add_instance_library INSTANCE_NAME)
+@@ -44,8 +44,13 @@ function(add_instance_library INSTANCE_NAME)
      endforeach()
      #only continue if there are some source files left on the list
      if(ARGN)
@@ -84,9 +77,13 @@ index 9cb5d0e9a..141a46f3d 100644
 +        # Always disable debug symbol and C debug assert due to
 +        # - Linker error: ... relocation truncated to fit ..., caused by object files to be linked are too huge.
 +        # - https://github.com/ROCmSoftwarePlatform/composable_kernel/issues/622
-+        target_compile_options(${INSTANCE_NAME} PRIVATE -g0 -DNDEBUG)
          target_compile_features(${INSTANCE_NAME} PUBLIC)
 +        target_compile_definitions(${INSTANCE_NAME} PRIVATE "__HIP_PLATFORM_AMD__=1" "__HIP_PLATFORM_HCC__=1")
          set_target_properties(${INSTANCE_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON)
          clang_tidy_check(${INSTANCE_NAME})
          set(result 0)
+@@ -280,4 +285,3 @@ set(DEV_OPS_INC_DIRS
+     ${PROJECT_SOURCE_DIR}/library/include/ck/
+ )
+ rocm_install(DIRECTORY ${DEV_OPS_INC_DIRS} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/ck)
+-