WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X

ROCm · Aug 24, 2023 · 6477e68 · 6477e68
1 parent 4d1b82a
commit 6477e68
Show file tree

Hide file tree

Showing 2 changed files with 26 additions and 3 deletions.
diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp
@@ -61,6 +61,10 @@
 #define MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_30 \
     (MIOPEN_USE_COMGR && BUILD_SHARED_LIBS && (HIP_PACKAGE_VERSION_FLAT < 4003000000ULL))
 
+/// hipMemGetInfo constantly fails on gfx906/900 and Navi21.
+/// Brute-force W/A: return fixed values.
+#define WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X (ROCM_FEATURE_DEPRECATED_VEGA_NAVI2X)
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEVICE_CU)
 
 namespace miopen {
@@ -72,10 +76,26 @@ void toCallHipInit() __attribute__((constructor(1000)));
 void toCallHipInit() { hipInit(0); }
 #endif
 
+hipError_t hip_mem_get_info_wrapper(std::size_t* const free, std::size_t* const total)
+{
+#if WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X
+    const auto status = hipMemGetInfo(free, total);
+    if(status == hipSuccess)
+        return status;
+    MIOPEN_LOG_W("hipMemGetInfo error, status: " << status);
+    assert(free != nullptr && total != nullptr);
+    *free  = 16ULL * 1024 * 1024 * 1024; // 16 GiB
+    *total = *free;
+    return hipSuccess;
+#else
+    return hipMemGetInfo(free, total);
+#endif
+}
+
 std::size_t GetAvailableMemory()
 {
     size_t free, total;
-    auto status = hipMemGetInfo(&free, &total);
+    auto status = hip_mem_get_info_wrapper(&free, &total);
     if(status != hipSuccess)
         MIOPEN_THROW_HIP_STATUS(status, "Failed getting available memory");
     return free;
@@ -289,7 +309,7 @@ Handle::Handle() : impl(std::make_unique<HandleImpl>())
     this->impl->device      = set_default_device();
     this->impl->root_stream = impl->create_stream();
 #else
-    this->impl->device = get_device_id();
+    this->impl->device      = get_device_id();
     this->impl->root_stream = HandleImpl::reference_stream(nullptr);
 #endif
     auto root_stream = this->impl->root_stream.get();
@@ -640,7 +660,7 @@ std::size_t Handle::GetMaxMemoryAllocSize()
     if(m_MaxMemoryAllocSizeCached == 0)
     {
         size_t free, total;
-        auto status = hipMemGetInfo(&free, &total);
+        auto status = hip_mem_get_info_wrapper(&free, &total);
         if(status != hipSuccess)
             MIOPEN_THROW_HIP_STATUS(status, "Failed getting available memory");
         m_MaxMemoryAllocSizeCached = floor(total * 0.85);

diff --git a/src/include/miopen/rocm_features.hpp b/src/include/miopen/rocm_features.hpp
@@ -77,4 +77,7 @@
 #define ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT \
     (HIP_PACKAGE_VERSION_FLAT >= 4001021072ULL)
 
+/// GFX906 and GFX103X are deprecated since 5.7 RC.
+#define ROCM_FEATURE_DEPRECATED_VEGA_NAVI2X (HIP_PACKAGE_VERSION_FLAT >= 5007000000ULL)
+
 #endif // GUARD_ROCM_FEATURES_HPP_