[ROCm] fix: obtain AMD GPU memory info through rocm_smi library (microsoft#21190)

### Description
Previously ROCMExecutionProvider uses `hipMemGetInfo` to obtain the
sizes of total memory and available memory. However, this API has been
broken since ROCm 5.7. In this PR, we use `rocm_smi` library instead of
`hipMemGetInfo`.


### Motivation and Context

`hipMemGetInfo` API has been broken since ROCm 5.7 and inference with
ROCMExecutionProvider will lead to following errors:

```
HIP failure 1: invalid argument ; GPU=0 ; hostname=4cc4900475fe ; file=/onnxruntime/onnxruntime/core/providers/rocm/rocm_execution_provider.cc ; line=229 ; expr=hipMemGetInfo(&free, &total);
```

MIOpen has a brute-force fix for this
(https://github.com/ROCm/MIOpen/blob/911e67189592c311374940493f2099f3abced60d/src/hip/handlehip.cpp#L72).
Instead of hard-coding available memory to 16GB, I suppose we could
obtain memory info through `rocm_smi` library as in this PR.

Author: hann-wang

cmake/onnxruntime_providers_rocm.cmake

@@ -49,7 +49,10 @@
 
   find_library(RCCL_LIB rccl REQUIRED)
   find_library(ROCTRACER_LIB roctracer64 REQUIRED)
-  set(ONNXRUNTIME_ROCM_LIBS roc::rocblas MIOpen hip::hipfft ${RCCL_LIB} ${ROCTRACER_LIB})
+  find_package(rocm_smi REQUIRED)
+  set(ONNXRUNTIME_ROCM_LIBS roc::rocblas MIOpen hip::hipfft ${ROCM_SMI_LIBRARY} ${RCCL_LIB} ${ROCTRACER_LIB})
+  include_directories(${ROCM_SMI_INCLUDE_DIR})
+  link_directories(${ROCM_SMI_LIB_DIR})
 
   file(GLOB_RECURSE onnxruntime_providers_rocm_cc_srcs CONFIGURE_DEPENDS
     "${ONNXRUNTIME_ROOT}/core/providers/rocm/*.h"

onnxruntime/core/providers/rocm/nn/conv.cc

@@ -49,10 +49,10 @@ miopenStatus_t GetWorkspaceSize(miopenHandle_t handle, const MiopenConvState<mio
 }
 
 size_t GetMaxWorkspaceSize(miopenHandle_t handle, const MiopenConvState<miopenConvAlgoPerf_t>& s,
-                           const miopenConvFwdAlgorithm_t* algo, int n_algo) {
+                           const miopenConvFwdAlgorithm_t* algo, int n_algo, int device_id = 0) {
   // TODO: get maximum available size from memory arena
   size_t free, total;
-  HIP_CALL_THROW(hipMemGetInfo(&free, &total));
+  onnxruntime::rocm::hipMemGetInfoAlt(device_id, &free, &total);
   // Assuming 10% of fragmentation
   free = static_cast<size_t>(static_cast<double>(free) * 0.9);
   size_t max_ws_size = 0;
@@ -283,7 +283,7 @@ Status Conv<T, NHWC>::UpdateState(OpKernelContext* context, bool bias_expected)
       int algo_count = 1;
       const ROCMExecutionProvider* rocm_ep = static_cast<const ROCMExecutionProvider*>(this->Info().GetExecutionProvider());
       static constexpr int num_algos = MIOPEN_CONVOLUTION_FWD_ALGO_COUNT;
-      size_t max_ws_size = rocm_ep->GetMiopenConvUseMaxWorkspace() ? GetMaxWorkspaceSize(GetMiopenHandle(context), s_, kAllAlgos, num_algos)
+      size_t max_ws_size = rocm_ep->GetMiopenConvUseMaxWorkspace() ? GetMaxWorkspaceSize(GetMiopenHandle(context), s_, kAllAlgos, num_algos, rocm_ep->GetDeviceId())
                                                                    : AlgoSearchWorkspaceSize;
       IAllocatorUniquePtr<void> algo_search_workspace = GetTransientScratchBuffer<void>(max_ws_size);
       MIOPEN_RETURN_IF_ERROR(miopenFindConvolutionForwardAlgorithm(

onnxruntime/core/providers/rocm/rocm_call.cc

@@ -143,6 +143,8 @@ template Status RocmCall<hiprandStatus_t, false>(hiprandStatus_t retCode, const
 template void RocmCall<hiprandStatus_t, true>(hiprandStatus_t retCode, const char* exprString, const char* libName, hiprandStatus_t successCode, const char* msg, const char* file, const int line);
 template Status RocmCall<hipfftResult, false>(hipfftResult retCode, const char* exprString, const char* libName, hipfftResult successCode, const char* msg, const char* file, const int line);
 template void RocmCall<hipfftResult, true>(hipfftResult retCode, const char* exprString, const char* libName, hipfftResult successCode, const char* msg, const char* file, const int line);
+template Status RocmCall<rsmi_status_t, false>(rsmi_status_t retCode, const char* exprString, const char* libName, rsmi_status_t successCode, const char* msg, const char* file, const int line);
+template void RocmCall<rsmi_status_t, true>(rsmi_status_t retCode, const char* exprString, const char* libName, rsmi_status_t successCode, const char* msg, const char* file, const int line);
 
 #ifdef ORT_USE_NCCL
 template Status RocmCall<ncclResult_t, false>(ncclResult_t retCode, const char* exprString, const char* libName, ncclResult_t successCode, const char* msg, const char* file, const int line);

onnxruntime/core/providers/rocm/rocm_common.h

@@ -67,5 +67,17 @@ inline int warpSizeDynamic() {
   return deviceProp.warpSize;
 }
 
+inline void hipMemGetInfoAlt(uint32_t deviceId, size_t* pFree, size_t* pTotal) {
+  const auto status = hipMemGetInfo(pFree, pTotal);
+  if (status != hipSuccess) {
+    size_t usedMemory = 0;
+    ROCMSMI_CALL_THROW(rsmi_init(0));
+    ROCMSMI_CALL_THROW(rsmi_dev_memory_total_get(deviceId, RSMI_MEM_TYPE_VIS_VRAM, pTotal));
+    ROCMSMI_CALL_THROW(rsmi_dev_memory_usage_get(deviceId, RSMI_MEM_TYPE_VIS_VRAM, &usedMemory));
+    *pFree = *pTotal - usedMemory;
+    ROCMSMI_CALL_THROW(rsmi_shut_down());
+  }
+}
+
 }  // namespace rocm
 }  // namespace onnxruntime

onnxruntime/core/providers/rocm/rocm_execution_provider.cc

@@ -261,7 +261,7 @@ ROCMExecutionProvider::ROCMExecutionProvider(const ROCMExecutionProviderInfo& in
 
   size_t free = 0;
   size_t total = 0;
-  HIP_CALL_THROW(hipMemGetInfo(&free, &total));
+  onnxruntime::rocm::hipMemGetInfoAlt(info_.device_id, &free, &total);
 
   OverrideTunableOpInfoByEnv(info_);
 

onnxruntime/core/providers/rocm/rocm_pch.h

@@ -14,6 +14,7 @@
 #include <hipsparse/hipsparse.h>
 #include <miopen/miopen.h>
 #include <rocblas/rocblas.h>
+#include <rocm_smi/rocm_smi.h>
 
 #ifdef ORT_USE_NCCL
 #include <rccl/rccl.h>

onnxruntime/core/providers/rocm/shared_inc/rocm_call.h

@@ -17,6 +17,7 @@ std::conditional_t<THRW, void, Status> RocmCall(
 
 #define HIP_CALL(expr) (RocmCall<hipError_t, false>((expr), #expr, "HIP", hipSuccess, "", __FILE__, __LINE__))
 #define ROCBLAS_CALL(expr) (RocmCall<rocblas_status, false>((expr), #expr, "ROCBLAS", rocblas_status_success, "", __FILE__, __LINE__))
+#define ROCMSMI_CALL(expr) (RocmCall<rsmi_status_t, false>((expr), #expr, "ROCMSMI", RSMI_STATUS_SUCCESS, "", __FILE__, __LINE__))
 
 #define HIPSPARSE_CALL(expr) (RocmCall<hipsparseStatus_t, false>((expr), #expr, "HIPSPARSE", HIPSPARSE_STATUS_SUCCESS, "", __FILE__, __LINE__))
 #define HIPRAND_CALL(expr) (RocmCall<hiprandStatus_t, false>((expr), #expr, "HIPRAND", HIPRAND_STATUS_SUCCESS, "", __FILE__, __LINE__))
@@ -27,6 +28,7 @@ std::conditional_t<THRW, void, Status> RocmCall(
 
 #define HIP_CALL_THROW(expr) (RocmCall<hipError_t, true>((expr), #expr, "HIP", hipSuccess, "", __FILE__, __LINE__))
 #define ROCBLAS_CALL_THROW(expr) (RocmCall<rocblas_status, true>((expr), #expr, "ROCBLAS", rocblas_status_success, "", __FILE__, __LINE__))
+#define ROCMSMI_CALL_THROW(expr) (RocmCall<rsmi_status_t, true>((expr), #expr, "ROCMSMI", RSMI_STATUS_SUCCESS, "", __FILE__, __LINE__))
 
 #define HIPSPARSE_CALL_THROW(expr) (RocmCall<hipsparseStatus_t, true>((expr), #expr, "HIPSPARSE", HIPSPARSE_STATUS_SUCCESS, "", __FILE__, __LINE__))
 #define HIPRAND_CALL_THROW(expr) (RocmCall<hiprandStatus_t, true>((expr), #expr, "HIPRAND", HIPRAND_STATUS_SUCCESS, "", __FILE__, __LINE__))