saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-12 17:57:38 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Update CK and fix performance due to lacking -amdgpu-early-inline-all=true (#13493)

Browse source 
1. Update CK to its latest develop branch
2. `-mllvm -amdgpu-early-inline-all=true` is critical to CK's
performance, add it.
This commit is contained in:
cloudhan 2022-10-29 00:36:00 +08:00 committed by GitHub
parent 8b0669bf63
commit 4dd053cc15
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
2 changed files with 11 additions and 16 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
cmake/CMakeLists.txt
View file

@ -1887,6 +1887,11 @@ if (onnxruntime_USE_ROCM)
set(CMAKE_HIP_ARCHITECTURES "gfx906;gfx908;gfx90a;gfx1030")
endif()
# NOTE: Following flags are extracted from hipcc via HIPCC_VERBOSE=7 hipcc foo.cc
# -mllvm pass the next flag to underlying llvm, not clang itself
# -amdgpu-early-inline-all=true is critical for kernel code
set(CMAKE_HIP_FLAGS "-mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false ${CMAKE_HIP_FLAGS}")
if (NOT onnxruntime_HIPIFY_PERL)
set(onnxruntime_HIPIFY_PERL ${onnxruntime_ROCM_HOME}/hip/bin/hipify-perl)
endif()

22
cmake/patches/composable_kernel/Fix_Clang_Build.patch
View file

@ -1,5 +1,5 @@
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 3e1174ec..65648cb7 100644
index 5655ba17..c6a10fe0 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,7 +1,7 @@
@ -11,7 +11,7 @@ index 3e1174ec..65648cb7 100644
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
@@ -36,27 +36,6 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON)
@@ -41,27 +41,6 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CXX_EXTENSIONS OFF)
message("CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")
@ -39,7 +39,7 @@ index 3e1174ec..65648cb7 100644
## HIP
find_package(HIP REQUIRED)
# Override HIP version in config.h, if necessary.
@@ -245,9 +224,6 @@ rocm_package_setup_component(tests
@@ -263,9 +242,6 @@ rocm_package_setup_component(tests
)
add_subdirectory(library)
@ -49,7 +49,7 @@ index 3e1174ec..65648cb7 100644
#Create an interface target for the include only files and call it "composablekernels"
include(CMakePackageConfigHelpers)
@@ -273,11 +249,3 @@ rocm_install(FILES
@@ -291,11 +267,3 @@ rocm_install(FILES
set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE")
set(CPACK_RPM_PACKAGE_LICENSE "MIT")
@ -62,7 +62,7 @@ index 3e1174ec..65648cb7 100644
- HEADER_ONLY
-)
diff --git a/include/ck/ck.hpp b/include/ck/ck.hpp
index fcaec592..8ea06421 100644
index 92018aac..2ada620c 100644
--- a/include/ck/ck.hpp
+++ b/include/ck/ck.hpp
@@ -126,7 +126,9 @@
@ -76,7 +76,7 @@ index fcaec592..8ea06421 100644
// hack: have underlying assumption that need to be satsified, otherwise it's a bug
diff --git a/library/src/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/CMakeLists.txt
index 6f3f900b..594d983d 100644
index c206c4dc..6ca1ec11 100644
--- a/library/src/tensor_operation_instance/gpu/CMakeLists.txt
+++ b/library/src/tensor_operation_instance/gpu/CMakeLists.txt
@@ -1,5 +1,6 @@
@ -86,13 +86,3 @@ index 6f3f900b..594d983d 100644
add_library(${INSTANCE_NAME} OBJECT ${ARGN})
target_compile_features(${INSTANCE_NAME} PUBLIC)
set_target_properties(${INSTANCE_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON)
diff --git a/library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
index 5dc20332..78eedca5 100644
--- a/library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
+++ b/library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt
@@ -1,4 +1,4 @@
-add_library(device_grouped_conv3d_fwd_instance
+add_instance_library(device_grouped_conv3d_fwd_instance
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_bf16_instance.cpp
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_f16_instance.cpp
device_grouped_conv3d_fwd_xdl_gndhwc_gkzyxc_gndhwk_f32_instance.cpp