3x faster MIOpen build by using host compiler instead of ROCm/clang++

CMakeLists.txt

@@ -23,7 +23,7 @@
 # SOFTWARE.
 #
 ################################################################################
-cmake_minimum_required( VERSION 3.11 )
+cmake_minimum_required( VERSION 3.24 FATAL_ERROR )
 
 macro(set_var_to_condition var)
     if(${ARGN})
@@ -46,7 +46,7 @@ else()
     set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "")
 endif()
 
-project ( MIOpen C CXX )
+project ( MIOpen C CXX HIP )
 
 enable_testing()
 
@@ -90,7 +90,6 @@ endif()
 rocm_setup_version(VERSION 2.19.0)
 
 list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
-include(TargetFlags)
 
 if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
     if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS "5.3")
@@ -191,7 +190,6 @@ endif()
 # HIP is always required
 find_package(hip REQUIRED PATHS /opt/rocm)
 message(STATUS "Build with HIP ${hip_VERSION}")
-target_flags(HIP_COMPILER_FLAGS hip::device)
 # Remove cuda arch flags
 string(REGEX REPLACE --cuda-gpu-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
 string(REGEX REPLACE --offload-arch=[a-z0-9:+-]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
@@ -242,7 +240,7 @@ option(MIOPEN_USE_HIPRTC "Use HIPRTC to build HIP kernels instead of COMGR" ${MI
 
 message(STATUS "Hip compiler flags: ${HIP_COMPILER_FLAGS}")
 
-add_definitions("-DHIP_COMPILER_FLAGS=${HIP_COMPILER_FLAGS}")
+#add_definitions("-DHIP_COMPILER_FLAGS=${HIP_COMPILER_FLAGS}")
 
 
 # HIP
@@ -271,10 +269,11 @@ if( MIOPEN_BACKEND STREQUAL "HIP" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_B
         message(STATUS "OpenCL compiler not found")
     endif()
 
-        # Hcc's clang always defines __HCC__ even when not using hcc driver
-        add_definitions(-U__HCC__)
+    # Hcc's clang always defines __HCC__ even when not using hcc driver
+    add_definitions(-U__HCC__)
 
-    set(MIOPEN_HIP_COMPILER ${CMAKE_CXX_COMPILER} CACHE PATH "")
+    enable_language(HIP)
+    set(MIOPEN_HIP_COMPILER ${CMAKE_HIP_COMPILER} CACHE PATH "")
 
     # rocblas
     set(MIOPEN_USE_ROCBLAS ON CACHE BOOL "")

README.md

@@ -103,7 +103,7 @@ cmake -DMIOPEN_BACKEND=HIP -DCMAKE_PREFIX_PATH="<hip-installed-path>;<rocm-insta
 An example cmake step can be:
 ```
 export CXX=/opt/rocm/llvm/bin/clang++ && \
-cmake -DMIOPEN_BACKEND=HIP -DCMAKE_PREFIX_PATH="/opt/rocm/;/opt/rocm/hip;/root/MIOpen/install_dir" ..
+cmake -DMIOPEN_BACKEND=HIP -DCMAKE_PREFIX_PATH="/opt/rocm/;/opt/rocm/hip;/root/MIOpen/install_dir" -DCMAKE_HIP_ARCHITECTURES="gfx900;gfx906;gfx908;gfx90a;gfx1030" ..
 ```
 
 Note: When specifying the path for the `CMAKE_PREFIX_PATH` variable, **do not** use the `~` shorthand for the user home directory.

include/miopen/config.h.in

@@ -154,4 +154,34 @@
 
 #define MIOPEN_GOLDEN_DB_VERSION 12
 
+// Otherwise, pArgs() is not defined, if compiled with HIP compiler.
+#ifndef __HIPCC__
+
+#include <cstddef>
+#include <tuple>
+#include <type_traits>
+
+template <std::size_t n, typename... Ts,
+          typename std::enable_if<n == sizeof...(Ts)>::type*>
+void pArgs(const std::tuple<Ts...>&, void*) {}
+
+template <std::size_t n, typename... Ts,
+          typename std::enable_if<n != sizeof...(Ts)>::type*>
+void pArgs(const std::tuple<Ts...>& formals, void** _vargs) {
+    using T = typename std::tuple_element<n, std::tuple<Ts...> >::type;
+
+    static_assert(!std::is_reference<T>{},
+                  "A __global__ function cannot have a reference as one of its "
+                  "arguments.");
+#if defined(HIP_STRICT)
+    static_assert(std::is_trivially_copyable<T>{},
+                  "Only TriviallyCopyable types can be arguments to a __global__ "
+                  "function");
+#endif
+    _vargs[n] = const_cast<void*>(reinterpret_cast<const void*>(&std::get<n>(formals)));
+    return pArgs<n + 1>(formals, _vargs);
+}
+
+#endif
+
 #endif // GUARD_CONFIG_H_IN

src/CMakeLists.txt

@@ -179,21 +179,21 @@ set( MIOpen_Source
     solver/conv_direct_naive_conv_bwd.cpp
     solver/conv_direct_naive_conv_fwd.cpp
     solver/conv_direct_naive_conv_wrw.cpp
-    solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp
-    solver/conv_hip_implicit_gemm_bwd_v1r1.cpp
-    solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp
-    solver/conv_hip_implicit_gemm_bwd_v4r1.cpp
-    solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp
-    solver/conv_hip_implicit_gemm_fwd_v4r1.cpp
-    solver/conv_hip_implicit_gemm_fwd_v4r4.cpp
-    solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp
-    solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp
-    solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp
-    solver/conv_hip_implicit_gemm_fwd_xdlops.cpp
-    solver/conv_hip_implicit_gemm_nonxdlops_common.cpp
-    solver/conv_hip_implicit_gemm_wrw_v4r4.cpp
-    solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp
-    solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp
+    solver/conv_hip_implicit_gemm_bwd_data_xdlops.hip
+    solver/conv_hip_implicit_gemm_bwd_v1r1.hip
+    solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.hip
+    solver/conv_hip_implicit_gemm_bwd_v4r1.hip
+    solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.hip
+    solver/conv_hip_implicit_gemm_fwd_v4r1.hip
+    solver/conv_hip_implicit_gemm_fwd_v4r4.hip
+    solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.hip
+    solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.hip
+    solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.hip
+    solver/conv_hip_implicit_gemm_fwd_xdlops.hip
+    solver/conv_hip_implicit_gemm_nonxdlops_common.hip
+    solver/conv_hip_implicit_gemm_wrw_v4r4.hip
+    solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.hip
+    solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.hip
     solver/conv_mlir_igemm_bwd.cpp
     solver/conv_mlir_igemm_bwd_xdlops.cpp
     solver/conv_mlir_igemm_fwd.cpp
@@ -652,11 +652,14 @@ target_include_directories(MIOpen PUBLIC
 
 set(MIOPEN_CK_LINK_FLAGS)
 if(MIOPEN_USE_COMPOSABLEKERNEL)
-set(MIOPEN_CK_LINK_FLAGS composable_kernel::device_operations hip::host)
+set(MIOPEN_CK_LINK_FLAGS composable_kernel::device_operations) # hip::host)
 endif()
 
 target_include_directories(MIOpen SYSTEM PUBLIC $<BUILD_INTERFACE:${HALF_INCLUDE_DIR}>)
 target_include_directories(MIOpen SYSTEM PRIVATE ${BZIP2_INCLUDE_DIR})
+set_target_properties(MIOpen PROPERTIES
+	HIP_SEPARABLE_COMPILATION ON
+	HIP_RESOLVE_DEVICE_SYMBOLS ON)
 target_link_libraries(MIOpen PRIVATE ${CMAKE_THREAD_LIBS_INIT} ${BZIP2_LIBRARIES} ${MIOPEN_CK_LINK_FLAGS})
 generate_export_header(MIOpen
     EXPORT_FILE_NAME ${PROJECT_BINARY_DIR}/include/miopen/export.h
@@ -671,14 +674,14 @@ if( MIOPEN_BACKEND STREQUAL "OpenCL")
     target_link_libraries( MIOpen PUBLIC ${OPENCL_LIBRARIES} )
     list(APPEND PACKAGE_DEPENDS PACKAGE OpenCL)
 elseif(MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_BACKEND STREQUAL "HIP")
-    target_link_libraries( MIOpen PRIVATE hip::device )
-    target_link_libraries( MIOpen INTERFACE hip::host )
+    #target_link_libraries( MIOpen PRIVATE hip::device )
+    #target_link_libraries( MIOpen INTERFACE hip::host )
     if(ENABLE_HIP_WORKAROUNDS)
         # Workaround hip not setting its usage requirements correctly
         target_compile_definitions( MIOpen PRIVATE -D__HIP_PLATFORM_HCC__=1 )
     endif()
     # This is helpful for the tests
-    target_link_libraries( MIOpen INTERFACE $<BUILD_INTERFACE:hip::device> )
+    #target_link_libraries( MIOpen INTERFACE $<BUILD_INTERFACE:hip::device> )
     list(APPEND PACKAGE_DEPENDS PACKAGE hip)
 endif()