Merge branch 'rocm6.2_internal_testing' into fix_cmake_migx_ep

cmake/CMakeLists.txt

@@ -1477,9 +1477,6 @@ if (onnxruntime_USE_TENSORRT)
 endif()
 
 if (onnxruntime_USE_MIGRAPHX)
-  if (WIN32)
-    message(FATAL_ERROR "MIGraphX does not support build in Windows!")
-  endif()
   set(AMD_MIGRAPHX_HOME ${onnxruntime_MIGRAPHX_HOME})
 endif()
 

cmake/onnxruntime_providers_migraphx.cmake

@@ -19,23 +19,25 @@
   endif()
 
   # Add search paths for default rocm installation
-  list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hcc /opt/rocm/hip /opt/rocm)
+  list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hcc /opt/rocm/hip /opt/rocm $ENV{HIP_PATH})
 
   # Suppress the warning about the small capitals of the package name - Enable when support to CMake 3.27.0 is used
   # cmake_policy(SET CMP0144 NEW)
 
-  find_package(miopen)
-  find_package(rocblas)
+  if(WIN32 AND NOT HIP_PLATFORM)
+    set(HIP_PLATFORM "amd")
+  endif()
+
+  find_package(hip REQUIRED)
+  find_package(migraphx REQUIRED PATHS ${AMD_MIGRAPHX_HOME})
 
-  set(migraphx_libs migraphx::c hip::host MIOpen roc::rocblas)
+  set(migraphx_libs migraphx::c hip::host)
 
   file(GLOB_RECURSE onnxruntime_providers_migraphx_cc_srcs CONFIGURE_DEPENDS
     "${ONNXRUNTIME_ROOT}/core/providers/migraphx/*.h"
     "${ONNXRUNTIME_ROOT}/core/providers/migraphx/*.cc"
     "${ONNXRUNTIME_ROOT}/core/providers/shared_library/*.h"
     "${ONNXRUNTIME_ROOT}/core/providers/shared_library/*.cc"
-    "${ONNXRUNTIME_ROOT}/core/providers/rocm/rocm_stream_handle.h"
-    "${ONNXRUNTIME_ROOT}/core/providers/rocm/rocm_stream_handle.cc"
   )
   source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_migraphx_cc_srcs})
   onnxruntime_add_shared_library_module(onnxruntime_providers_migraphx ${onnxruntime_providers_migraphx_cc_srcs})
@@ -46,18 +48,16 @@
   set_target_properties(onnxruntime_providers_migraphx PROPERTIES LINKER_LANGUAGE CXX)
   set_target_properties(onnxruntime_providers_migraphx PROPERTIES FOLDER "ONNXRuntime")
   target_compile_definitions(onnxruntime_providers_migraphx PRIVATE ONNXIFI_BUILD_LIBRARY=1)
-  target_compile_options(onnxruntime_providers_migraphx PRIVATE -Wno-error=sign-compare)
-  set_property(TARGET onnxruntime_providers_migraphx APPEND_STRING PROPERTY COMPILE_FLAGS "-Wno-deprecated-declarations")
-  set_property(TARGET onnxruntime_providers_migraphx APPEND_STRING PROPERTY LINK_FLAGS "-Xlinker --version-script=${ONNXRUNTIME_ROOT}/core/providers/migraphx/version_script.lds -Xlinker --gc-sections")
-  target_link_libraries(onnxruntime_providers_migraphx PRIVATE nsync::nsync_cpp stdc++fs)
-
-  include(CheckLibraryExists)
-  check_library_exists(migraphx::c "migraphx_program_run_async" "/opt/rocm/migraphx/lib" HAS_STREAM_SYNC)
-  if(HAS_STREAM_SYNC)
-      target_compile_definitions(onnxruntime_providers_migraphx PRIVATE -DMIGRAPHX_STREAM_SYNC)
-      message(STATUS "MIGRAPHX GPU STREAM SYNC is ENABLED")
+  if(MSVC)
+    set_property(TARGET onnxruntime_providers_migraphx APPEND_STRING PROPERTY LINK_FLAGS /DEF:${ONNXRUNTIME_ROOT}/core/providers/migraphx/symbols.def)
+    target_link_libraries(onnxruntime_providers_migraphx PRIVATE ws2_32)
   else()
-      message(STATUS "MIGRAPHX GPU STREAM SYNC is DISABLED")
+    target_compile_options(onnxruntime_providers_migraphx PRIVATE -Wno-error=sign-compare)
+    set_property(TARGET onnxruntime_providers_migraphx APPEND_STRING PROPERTY COMPILE_FLAGS "-Wno-deprecated-declarations")
+  endif()
+  if(UNIX)
+    set_property(TARGET onnxruntime_providers_migraphx APPEND_STRING PROPERTY LINK_FLAGS "-Xlinker --version-script=${ONNXRUNTIME_ROOT}/core/providers/migraphx/version_script.lds -Xlinker --gc-sections")
+    target_link_libraries(onnxruntime_providers_migraphx PRIVATE nsync::nsync_cpp stdc++fs)
   endif()
 
   if (onnxruntime_ENABLE_TRAINING_OPS)
@@ -68,8 +68,16 @@
     endif()
   endif()
 
-  install(TARGETS onnxruntime_providers_migraphx
-          ARCHIVE  DESTINATION ${CMAKE_INSTALL_LIBDIR}
-          LIBRARY  DESTINATION ${CMAKE_INSTALL_LIBDIR}
-          RUNTIME  DESTINATION ${CMAKE_INSTALL_BINDIR}
-  )
+  if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
+    install(TARGETS onnxruntime_providers_migraphx
+            ARCHIVE  DESTINATION ${CMAKE_INSTALL_LIBDIR}
+            LIBRARY  DESTINATION ${CMAKE_INSTALL_BINDIR}
+            RUNTIME  DESTINATION ${CMAKE_INSTALL_BINDIR}
+    )
+  else()
+    install(TARGETS onnxruntime_providers_migraphx
+            ARCHIVE  DESTINATION ${CMAKE_INSTALL_LIBDIR}
+            LIBRARY  DESTINATION ${CMAKE_INSTALL_LIBDIR}
+            RUNTIME  DESTINATION ${CMAKE_INSTALL_BINDIR}
+    )
+  endif()

include/onnxruntime/core/session/onnxruntime_c_api.h

@@ -608,6 +608,10 @@ typedef struct OrtMIGraphXProviderOptions {
   int migraphx_int8_enable;                          // MIGraphX INT8 precision. Default 0 = false, nonzero = true
   int migraphx_use_native_calibration_table;         // MIGraphx INT8 cal table. Default 0 = false, noznero = true
   const char* migraphx_int8_calibration_table_name;  // MIGraphx INT8 calibration table name
+  int migraphx_save_compiled_model;                  // migraphx save compiled model. Default 0 = false, noznero = true
+  const char* migraphx_save_model_path;              // migraphx model path name
+  int migraphx_load_compiled_model;                  // migraphx int8 cal table. Default 0 = false, noznero = true
+  const char* migraphx_load_model_path;              // migraphx model path name
 } OrtMIGraphXProviderOptions;
 
 /** \brief OpenVINO Provider Options

onnxruntime/core/providers/migraphx/gpu_data_transfer.cc

@@ -60,17 +60,7 @@ common::Status GPUDataTransfer::CopyTensorAsync(const Tensor& src, Tensor& dst,
       HIP_CALL_THROW(hipMemcpy(dst_data, src_data, bytes, hipMemcpyHostToDevice));
     }
   } else if (src_device.Type() == OrtDevice::GPU) {
-#ifndef MIGRAPHX_STREAM_SYNC
-    if (dst_device.Type() == OrtDevice::CPU && dst_device.MemType() == OrtDevice::MemType::HIP_PINNED) {
-      // copying from GPU to pinned memory, this is non-blocking
-      HIP_CALL_THROW(hipMemcpyAsync(dst_data, src_data, bytes, hipMemcpyDeviceToHost, static_cast<hipStream_t>(stream.GetHandle())));
-    } else {
-      // copying from GPU to CPU memory, this is blocking
-      HIP_CALL_THROW(hipMemcpy(dst_data, src_data, bytes, hipMemcpyDeviceToHost));
-    }
-#else
     HIP_CALL_THROW(hipMemcpyAsync(dst_data, src_data, bytes, hipMemcpyDeviceToHost, static_cast<hipStream_t>(stream.GetHandle())));
-#endif
   } else {
     // copying between cpu memory
     memcpy(dst_data, src_data, bytes);

onnxruntime/core/providers/migraphx/hip_allocator.cc → onnxruntime/core/providers/migraphx/migraphx_allocator.cc

@@ -3,15 +3,15 @@
 
 #include "core/providers/shared_library/provider_api.h"
 #include "migraphx_call.h"
-#include "hip_allocator.h"
+#include "migraphx_allocator.h"
 #include "core/common/status.h"
 #include "core/framework/float16.h"
 #include "core/common/status.h"
 #include "gpu_data_transfer.h"
 
 namespace onnxruntime {
 
-void HIPAllocator::CheckDevice() const {
+void MIGraphXAllocator::CheckDevice() const {
 #ifndef NDEBUG
   // check device to match at debug build
   // if it's expected to change, call hipSetDevice instead of the check
@@ -23,7 +23,7 @@ void HIPAllocator::CheckDevice() const {
 #endif
 }
 
-void* HIPAllocator::Alloc(size_t size) {
+void* MIGraphXAllocator::Alloc(size_t size) {
   CheckDevice();
   void* p = nullptr;
   if (size > 0) {
@@ -32,12 +32,12 @@ void* HIPAllocator::Alloc(size_t size) {
   return p;
 }
 
-void HIPAllocator::Free(void* p) {
+void MIGraphXAllocator::Free(void* p) {
   CheckDevice();
   (void)hipFree(p);  // do not throw error since it's OK for hipFree to fail during shutdown
 }
 
-void* HIPExternalAllocator::Alloc(size_t size) {
+void* MIGraphXExternalAllocator::Alloc(size_t size) {
   void* p = nullptr;
   if (size > 0) {
     p = alloc_(size);
@@ -49,7 +49,7 @@ void* HIPExternalAllocator::Alloc(size_t size) {
   return p;
 }
 
-void HIPExternalAllocator::Free(void* p) {
+void MIGraphXExternalAllocator::Free(void* p) {
   free_(p);
   std::lock_guard<OrtMutex> lock(lock_);
   auto it = reserved_.find(p);
@@ -59,7 +59,7 @@ void HIPExternalAllocator::Free(void* p) {
   }
 }
 
-void* HIPExternalAllocator::Reserve(size_t size) {
+void* MIGraphXExternalAllocator::Reserve(size_t size) {
   void* p = Alloc(size);
   if (!p) return nullptr;
   std::lock_guard<OrtMutex> lock(lock_);

onnxruntime/core/providers/migraphx/hip_allocator.h → onnxruntime/core/providers/migraphx/migraphx_allocator.h

@@ -9,12 +9,12 @@
 
 namespace onnxruntime {
 
-class HIPAllocator : public IAllocator {
+class MIGraphXAllocator : public IAllocator {
  public:
-  HIPAllocator(int device_id, const char* name)
+  MIGraphXAllocator(int device_id, const char* name)
       : IAllocator(
             OrtMemoryInfo(name, OrtAllocatorType::OrtDeviceAllocator,
-                          OrtDevice(OrtDevice::GPU, OrtDevice::MemType::DEFAULT, device_id),
+                          OrtDevice(OrtDevice::GPU, OrtDevice::MemType::DEFAULT, static_cast<OrtDevice::DeviceId>(device_id)),
                           device_id, OrtMemTypeDefault)) {}
 
   virtual void* Alloc(size_t size) override;
@@ -24,14 +24,14 @@ class HIPAllocator : public IAllocator {
   void CheckDevice() const;
 };
 
-class HIPExternalAllocator : public HIPAllocator {
+class MIGraphXExternalAllocator : public MIGraphXAllocator {
   typedef void* (*ExternalAlloc)(size_t size);
   typedef void (*ExternalFree)(void* p);
   typedef void (*ExternalEmptyCache)();
 
  public:
-  HIPExternalAllocator(OrtDevice::DeviceId device_id, const char* name, void* alloc, void* free, void* empty_cache)
-      : HIPAllocator(device_id, name) {
+  MIGraphXExternalAllocator(OrtDevice::DeviceId device_id, const char* name, void* alloc, void* free, void* empty_cache)
+      : MIGraphXAllocator(device_id, name) {
     alloc_ = reinterpret_cast<ExternalAlloc>(alloc);
     free_ = reinterpret_cast<ExternalFree>(free);
     empty_cache_ = reinterpret_cast<ExternalEmptyCache>(empty_cache);
@@ -55,7 +55,7 @@ class HIPPinnedAllocator : public IAllocator {
   HIPPinnedAllocator(int device_id, const char* name)
       : IAllocator(
             OrtMemoryInfo(name, OrtAllocatorType::OrtDeviceAllocator,
-                          OrtDevice(OrtDevice::CPU, OrtDevice::MemType::HIP_PINNED, device_id),
+                          OrtDevice(OrtDevice::CPU, OrtDevice::MemType::HIP_PINNED, static_cast<OrtDevice::DeviceId>(device_id)),
                           device_id, OrtMemTypeCPUOutput)) {}
 
   virtual void* Alloc(size_t size) override;

onnxruntime/core/providers/migraphx/migraphx_call.cc

@@ -1,10 +1,13 @@
 // Copyright (c) Microsoft Corporation. All rights reserved.
 // Licensed under the MIT License.
 
+#ifdef _WIN32
+#include <winsock.h>
+#else
 #include <unistd.h>
-#include <string.h>
-#include <miopen/miopen.h>
-#include <rocblas/rocblas.h>
+#endif
+
+#include <string>
 #include "core/common/common.h"
 #include "core/common/status.h"
 #include "core/providers/shared_library/provider_api.h"
@@ -34,16 +37,20 @@ std::conditional_t<THRW, void, Status> RocmCall(
     ERRTYPE retCode, const char* exprString, const char* libName, ERRTYPE successCode, const char* msg, const char* file, const int line) {
   if (retCode != successCode) {
     try {
-      char hostname[HOST_NAME_MAX];
-      if (gethostname(hostname, HOST_NAME_MAX) != 0)
-        strcpy(hostname, "?");
+#ifdef _WIN32
+      // According to the POSIX spec, 255 is the safe minimum value.
+      static constexpr int HOST_NAME_MAX = 255;
+#endif
+      std::string hostname(HOST_NAME_MAX, 0);
+      if (gethostname(hostname.data(), HOST_NAME_MAX) != 0)
+        hostname = "?";
       int currentHipDevice;
       (void)hipGetDevice(&currentHipDevice);
       (void)hipGetLastError();  // clear last HIP error
       static char str[1024];
       snprintf(str, 1024, "%s failure %d: %s ; GPU=%d ; hostname=%s ; file=%s ; line=%d ; expr=%s; %s",
                libName, (int)retCode, RocmErrString(retCode), currentHipDevice,
-               hostname,
+               hostname.c_str(),
                file, line, exprString, msg);
       if constexpr (THRW) {
         // throw an exception with the error info
@@ -68,9 +75,5 @@ std::conditional_t<THRW, void, Status> RocmCall(
 
 template Status RocmCall<hipError_t, false>(hipError_t retCode, const char* exprString, const char* libName, hipError_t successCode, const char* msg, const char* file, const int line);
 template void RocmCall<hipError_t, true>(hipError_t retCode, const char* exprString, const char* libName, hipError_t successCode, const char* msg, const char* file, const int line);
-template Status RocmCall<rocblas_status, false>(rocblas_status retCode, const char* exprString, const char* libName, rocblas_status successCode, const char* msg, const char* file, const int line);
-template void RocmCall<rocblas_status, true>(rocblas_status retCode, const char* exprString, const char* libName, rocblas_status successCode, const char* msg, const char* file, const int line);
-template Status RocmCall<miopenStatus_t, false>(miopenStatus_t retCode, const char* exprString, const char* libName, miopenStatus_t successCode, const char* msg, const char* file, const int line);
-template void RocmCall<miopenStatus_t, true>(miopenStatus_t retCode, const char* exprString, const char* libName, miopenStatus_t successCode, const char* msg, const char* file, const int line);
 
 }  // namespace onnxruntime

onnxruntime/core/providers/migraphx/migraphx_call.h

@@ -4,8 +4,6 @@
 #pragma once
 #include "migraphx_inc.h"
 
-#pragma once
-
 namespace onnxruntime {
 
 // -----------------------------------------------------------------------