Add remove MIOpen build option (ROCm#3124)

CMakeLists.txt

@@ -63,6 +63,12 @@ endif()
 
 option(MIGRAPHX_USE_ROCBLAS "Enable MIGraphX to use rocBLAS" ON)
 
+if(WIN32)
+option(MIGRAPHX_USE_MIOPEN "Enable MIGraphX to use MIOpen" OFF)
+else()
+option(MIGRAPHX_USE_MIOPEN "Enable MIGraphX to use MIOpen" ON)
+endif()
+
 # By default build shared libraries
 option(BUILD_SHARED_LIBS "Create shared libraries" ON)
 
@@ -300,6 +306,7 @@ rocm_enable_cppcheck(
         MIGRAPHX_MLIR=1
         MIGRAPHX_HAS_EXECUTORS=0
         CPPCHECK=1
+        MIGRAPHX_USE_MIOPEN=1
         BUILD_DEV=
         __device__=
         __host__=
@@ -339,6 +346,9 @@ else()
     set(DEPENDS_HIP_RUNTIME "hip-runtime-amd" )
 endif()
 
+if(MIGRAPHX_USE_MIOPEN)
+    list(APPEND PACKAGE_DEPENDS miopen-hip)
+endif()
 if(MIGRAPHX_USE_ROCBLAS)
     list(APPEND PACKAGE_DEPENDS rocblas)
 endif()
@@ -352,5 +362,5 @@ rocm_create_package(
     MAINTAINER "AMDMIGraphX Maintainer <[email protected]>"
     LDCONFIG
     PTH
-    DEPENDS miopen-hip ${DEPENDS_HIP_RUNTIME} half ${PACKAGE_DEPENDS}
+    DEPENDS ${DEPENDS_HIP_RUNTIME} half ${PACKAGE_DEPENDS}
 )

src/CMakeLists.txt

@@ -333,10 +333,12 @@ target_link_libraries(migraphx_all_targets INTERFACE migraphx_cpu)
 target_compile_definitions(migraphx_all_targets INTERFACE -DHAVE_CPU)
 endif()
 if(MIGRAPHX_ENABLE_GPU)
+    if(MIGRAPHX_USE_MIOPEN)
+        list(APPEND MIGRAPHX_CONFIG_DEPENDS PACKAGE MIOpen)
+    endif()
     if(MIGRAPHX_USE_ROCBLAS)
         list(APPEND MIGRAPHX_CONFIG_DEPENDS PACKAGE rocblas)
     endif()
-    list(APPEND MIGRAPHX_CONFIG_DEPENDS PACKAGE MIOpen)
 add_subdirectory(targets/gpu)
 target_link_libraries(migraphx_all_targets INTERFACE migraphx_gpu)
 target_compile_definitions(migraphx_all_targets INTERFACE -DHAVE_GPU)
@@ -347,6 +349,12 @@ target_link_libraries(migraphx_all_targets INTERFACE migraphx_fpga)
 target_compile_definitions(migraphx_all_targets INTERFACE -DHAVE_FPGA)
 endif()
 
+if(MIGRAPHX_USE_MIOPEN)
+    target_compile_definitions(migraphx_all_targets INTERFACE MIGRAPHX_USE_MIOPEN=1)
+else()
+    target_compile_definitions(migraphx_all_targets INTERFACE MIGRAPHX_USE_MIOPEN=0)
+endif()
+
 if(HAVE_HALF_EXPR)
     target_compile_definitions(migraphx PUBLIC -DHAS_HALF_V1)
 endif()

src/targets/gpu/CMakeLists.txt

@@ -30,8 +30,13 @@ if(NOT GPU_TARGETS)
     endif()
     message(FATAL_ERROR ${fatal_msg})
 endif()
-find_package(miopen REQUIRED)
-message(STATUS "MIGraphX is using MIOpen")
+
+if(MIGRAPHX_USE_MIOPEN)
+    find_package(miopen REQUIRED)
+    message(STATUS "MIGraphX is using MIOpen")
+else()
+    message(STATUS "MIGraphX is not using MIOpen")
+endif()
 
 if(MIGRAPHX_USE_ROCBLAS)
     # rocblas
@@ -85,6 +90,9 @@ rocm_set_soversion(migraphx_device ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_device)
 target_link_libraries(migraphx_device PUBLIC migraphx)
 target_link_libraries(migraphx_device PRIVATE compile_for_gpu)
+if(NOT MIGRAPHX_USE_MIOPEN AND NOT MIGRAPHX_USE_ROCBLAS)
+    target_link_libraries(migraphx_device INTERFACE hip::host)
+endif()
 target_include_directories(migraphx_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
 target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_BINAR_DIR}/include>)
 target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device/include>)
@@ -116,8 +124,11 @@ if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
             ${CMAKE_CURRENT_SOURCE_DIR}/jit/ck_gemm_softmax_gemm.cpp)
 endif()
 
+if(MIGRAPHX_USE_MIOPEN)
+    set(MIOPEN_SRCS abs.cpp)
+endif()
+
 add_library(migraphx_gpu
-    abs.cpp
     analyze_streams.cpp
     allocation_model.cpp
     argmax.cpp
@@ -161,6 +172,7 @@ add_library(migraphx_gpu
     topk.cpp
     write_literals.cpp
     ${JIT_GPU_SRCS}
+    ${MIOPEN_SRCS}
 )
 
 set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
@@ -183,12 +195,18 @@ register_migraphx_gpu_ops(hip_
     reverse
     topk
 )
+if (MIGRAPHX_USE_MIOPEN)
 register_migraphx_gpu_ops(miopen_
     abs
     contiguous
     lrn
     pooling
 )
+else()
+register_migraphx_gpu_ops(miopen_
+    contiguous
+)
+endif()
 register_op(migraphx_gpu
     HEADER migraphx/gpu/rnn_variable_seq_lens.hpp
     OPERATORS gpu::hip_rnn_var_sl_shift_sequence gpu::hip_rnn_var_sl_shift_output gpu::hip_rnn_var_sl_last_output
@@ -199,9 +217,11 @@ if(MIGRAPHX_USE_ROCBLAS)
         OPERATORS gpu::rocblas_gemm<op::dot> gpu::rocblas_gemm<op::quant_dot>
         INCLUDES migraphx/gpu/context.hpp)
 endif()
-register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
-    OPERATORS gpu::miopen_convolution<op::convolution> gpu::miopen_convolution<op::convolution_backwards> gpu::miopen_convolution<op::quant_convolution>
-    INCLUDES migraphx/gpu/context.hpp)
+if (MIGRAPHX_USE_MIOPEN)
+    register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
+        OPERATORS gpu::miopen_convolution<op::convolution> gpu::miopen_convolution<op::convolution_backwards> gpu::miopen_convolution<op::quant_convolution>
+        INCLUDES migraphx/gpu/context.hpp)
+endif()
 rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_gpu)
 
@@ -265,9 +285,14 @@ target_compile_definitions(migraphx_gpu PUBLIC MIGRAPHX_CXX_COMPILER="${CMAKE_CX
 # Check miopen find mode api
 
 include(CheckLibraryExists)
-get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
-check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
-check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
+if (MIGRAPHX_USE_MIOPEN)
+    get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
+    target_compile_definitions(migraphx_gpu PUBLIC MIGRAPHX_USE_MIOPEN=1)
+    check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
+    check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
+else()
+target_compile_definitions(migraphx_gpu PUBLIC MIGRAPHX_USE_MIOPEN=0)
+endif()
 
 if(MIGRAPHX_USE_ROCBLAS)
     get_target_property(ROCBLAS_LOCATION roc::rocblas LOCATION)
@@ -280,25 +305,29 @@ else()
     target_compile_definitions(migraphx_gpu PUBLIC MIGRAPHX_USE_ROCBLAS=0)
 endif()
 
-set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
+if(MIGRAPHX_USE_MIOPEN)
+    set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
 
-if(MIGRAPHX_USE_FIND_2_API)
-    check_library_exists(MIOpen "miopenSetFindOptionPreallocatedTensor" "${MIOPEN_LOCATION}" HAS_PREALLOCATION_API)
-    if(HAS_PREALLOCATION_API)
-        target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API -DMIGRAPHX_PREALLOCATE_MIOPEN_BUFFERS)
+    if(MIGRAPHX_USE_FIND_2_API)
+        check_library_exists(MIOpen "miopenSetFindOptionPreallocatedTensor" "${MIOPEN_LOCATION}" HAS_PREALLOCATION_API)
+        if(HAS_PREALLOCATION_API)
+            target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API -DMIGRAPHX_PREALLOCATE_MIOPEN_BUFFERS)
+        else()
+            target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API)
+        endif()
+        message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
     else()
-        target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API)
+        message(STATUS "MIGraphx is using legacy Find API in MIOpen")
     endif()
-    message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
-else()
-    message(STATUS "MIGraphx is using legacy Find API in MIOpen")
-endif()
 
-if(HAS_FIND_MODE_API)
-    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
-    message(STATUS "MIGraphx is using Find Mode API of MIOpen")
-else()
-    message(STATUS "MIOpen does not have find mode api")
+    if(HAS_FIND_MODE_API)
+        target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
+        message(STATUS "MIGraphx is using Find Mode API of MIOpen")
+    else()
+        message(STATUS "MIOpen does not have find mode api")
+    endif()
+
+    target_link_libraries(migraphx_gpu PUBLIC MIOpen)
 endif()
 
 if(MIGRAPHX_USE_ROCBLAS)
@@ -320,9 +349,13 @@ if(MIGRAPHX_USE_ROCBLAS)
     target_link_libraries(migraphx_gpu PUBLIC roc::rocblas)
 endif()
 
-target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen)
-
-target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
+target_link_libraries(migraphx_gpu PUBLIC migraphx)
+if(NOT MIGRAPHX_USE_MIOPEN AND NOT MIGRAPHX_USE_ROCBLAS)
+    target_link_libraries(migraphx_gpu PUBLIC migraphx_device)
+else()
+    target_link_libraries(migraphx_gpu PRIVATE migraphx_device)
+endif()
+target_link_libraries(migraphx_gpu PRIVATE migraphx_kernels)
 if(MIGRAPHX_USE_COMPOSABLEKERNEL)
     target_link_libraries(migraphx_gpu PRIVATE composable_kernel::jit_library)
     target_compile_definitions(migraphx_gpu PRIVATE MIGRAPHX_USE_COMPOSABLEKERNEL=1)

src/targets/gpu/abs.cpp

@@ -27,7 +27,7 @@
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
-
+#if MIGRAPHX_USE_MIOPEN
 shape miopen_abs::compute_shape(const std::vector<shape>& inputs) const
 {
     check_shapes{inputs, *this}.has(2).packed();
@@ -55,7 +55,7 @@ argument miopen_abs::compute(context& ctx,
 }
 
 void miopen_abs::finalize(context&, const shape&, const std::vector<shape>&) { ad = make_abs(); }
-
+#endif
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/fuse_ops.cpp

@@ -42,7 +42,7 @@ inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_MIOPEN_FUSION)
-
+#if MIGRAPHX_USE_MIOPEN
 struct fusion
 {
     using op_t = miopenFusionOpDescriptor_t;
@@ -162,14 +162,15 @@ struct fusion
         return y;
     }
 };
+#endif
 
 const std::unordered_set<std::string>& get_supported_archs()
 {
     static std::unordered_set<std::string> supported_archs{
         "gfx900", "gfx906", "gfx908", "gfx1030", "gfx940", "gfx941", "gfx942"};
     return supported_archs;
 }
-
+#if MIGRAPHX_USE_MIOPEN
 MIGRAPHX_PRED_MATCHER(bias_shape, instruction_ref ins)
 {
     auto&& s = ins->get_shape();
@@ -210,6 +211,7 @@ MIGRAPHX_PRED_MATCHER(fusable_conv, instruction_ref ins)
     return contains({{0, 0, 0, 0}, {1, 1, 1, 1}, {2, 2, 2, 2}}, conv_op.padding) and
            contains({{0, 0}, {1, 1}}, conv_op.stride) and contains({{1, 1}}, conv_op.dilation);
 }
+#endif
 
 void move_broadcasted_back(std::vector<instruction_ref>& args)
 {
@@ -234,6 +236,7 @@ void move_standard_front(std::vector<instruction_ref>& args)
 auto gpu_name(const std::string& s) { return match::name("gpu::" + s); }
 
 namespace {
+#if MIGRAPHX_USE_MIOPEN
 struct miopen_fusion
 {
     struct fuse_op_data
@@ -473,6 +476,7 @@ void apply_conv_bias(context& ctx, module& m, const match::matcher_result& r)
     (void)ws;
     m.replace_instruction(ins, cb, input_ins, weights_ins, old_ws_ins, bias_ins, alloc_ins);
 }
+#endif
 
 template <class... Strings>
 inline auto precompile_name(Strings... names) // NOLINT
@@ -485,6 +489,7 @@ inline auto precompile_name(Strings... names) // NOLINT
     });
 }
 
+#if MIGRAPHX_USE_MIOPEN
 struct find_conv_bias
 {
     context* ctx = nullptr;
@@ -510,7 +515,6 @@ struct find_conv_bias_relu
         apply_conv_bias<miopen_conv_bias_relu>(*ctx, m, r);
     }
 };
-
 struct find_conv_pointwise
 {
     context* ctx = nullptr;
@@ -549,6 +553,7 @@ struct find_conv_pointwise
         m.replace_instruction(ins, op, inputs);
     }
 };
+#endif
 
 #if MIGRAPHX_USE_ROCBLAS
 struct find_gemm_pointwise
@@ -892,8 +897,10 @@ void fuse_ops::apply(module& m) const
 {
     match::find_matches(m, find_pointwise_layout_contiguous{}, find_contiguous_layout_pointwise{});
     run_passes(m, {dead_code_elimination{}});
+#if MIGRAPHX_USE_MIOPEN
     match::find_matches(m, find_conv_pointwise{ctx}, find_conv_bias_relu{ctx}, find_conv_bias{ctx});
     run_passes(m, {dead_code_elimination{}});
+#endif
     match::find_matches(m,
 #if MIGRAPHX_USE_ROCBLAS
                         find_gemm_pointwise{},

src/targets/gpu/hip.cpp

@@ -27,7 +27,9 @@
 #include <migraphx/register_op.hpp>
 #include <migraphx/gpu/context.hpp>
 #include <migraphx/gpu/device/contiguous.hpp>
+#if MIGRAPHX_USE_MIOPEN
 #include <miopen/miopen.h>
+#endif
 #include <memory>
 #include <mutex>
 #include <vector>

src/targets/gpu/hiprtc/CMakeLists.txt

@@ -1,7 +1,7 @@
 #####################################################################################
 # The MIT License (MIT)
 #
-# Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal

src/targets/gpu/include/migraphx/gpu/abs.hpp

@@ -34,6 +34,7 @@ inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 struct context;
+#if MIGRAPHX_USE_MIOPEN
 
 struct miopen_abs
 {
@@ -56,7 +57,7 @@ struct miopen_abs
         return shapes.size() - 1;
     }
 };
-
+#endif
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx