From bab334057647bd85de390f07f59a0f7e7c9713ce Mon Sep 17 00:00:00 2001
From: Artem Tamazov <artem.tamazov@gmail.com>
Date: Wed, 14 Feb 2024 07:33:02 +0300
Subject: [PATCH] Fix for issue #2734: Detect if "-fno-offload-uniform-block"
 works in HIP compiler. (#2743)

* fix-issue-2734 (01) Use "-fno-offload-uniform-block" only if HIP compiler supports it. Resolves #2734.

(cherry picked from commit 458c8338175383a95a5c3f30c726798828f15ea8)

Partially changes code from PR #2719 "Do not use HIP runtime headers on Windows"

# RESOLVED Conflicts:
#	CMakeLists.txt

* fix-issue-2734(02) Removed W/A from PR #2719 as it is no longer needed.
---
 CMakeLists.txt             | 20 ++++++++++++++++++--
 include/miopen/config.h.in |  1 +
 src/comgr.cpp              | 12 +++---------
 3 files changed, 22 insertions(+), 11 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 32d9a2e5b0..76fa626656 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -258,8 +258,24 @@ set(MIOPEN_hip_VERSION ${MIOPEN_hip_VERSION_MAJOR}.${MIOPEN_hip_VERSION_MINOR}.$
 set_var_to_condition(MIOPEN_USE_HIPRTC_DEFAULT (MIOPEN_USE_COMGR AND (MIOPEN_hip_VERSION VERSION_GREATER_EQUAL 5)))
 option(MIOPEN_USE_HIPRTC "Use HIPRTC to build HIP kernels instead of COMGR" ${MIOPEN_USE_HIPRTC_DEFAULT})
 
-# WORKAROUND_SWDEV_413293 - do not use on Windows. Compiler error message: unknown command line option
-if(NOT WIN32 AND MIOPEN_hip_VERSION VERSION_GREATER_EQUAL 5.7.23302)
+# WORKAROUND_SWDEV_413293
+# Assume that any HIP kernel can be launched with non-uniform block size; otherwise
+# the "Failed to launch kernel: invalid argument" error may happen at run time.
+# References: SWDEV-413293 and https://reviews.llvm.org/D155213 effective HIP_FLAT_VERSION 500723302 on Linux.
+# This may lead to perf drops in the future therefore https://github.com/ROCm/MIOpen/issues/2708 is opened.
+if(HAS_HIP)
+    # HIP version is unreliable on Windows and on Fedora, so we use compiler flag detection,
+    # if this is possible. See issue 2734 and PR 2719.
+    check_cxx_compiler_flag("-x hip -fno-offload-uniform-block" MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK)
+else()
+    # CXX compiler is not HIP compiler, let's analyze HIP version.
+    set(MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK Off)
+    if(${MIOPEN_hip_VERSION_FLAT} GREATER_EQUAL 500723302)
+        set(MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK On)
+    endif()
+    message(STATUS "MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK: ${MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK}")
+endif()
+if(MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK)
     string(APPEND HIP_COMPILER_FLAGS " -fno-offload-uniform-block ")
 endif()
 
diff --git a/include/miopen/config.h.in b/include/miopen/config.h.in
index b66ce96380..6c1a867b59 100644
--- a/include/miopen/config.h.in
+++ b/include/miopen/config.h.in
@@ -50,6 +50,7 @@
 #cmakedefine01 MIOPEN_USE_COMPOSABLEKERNEL
 #cmakedefine01 MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK
 #cmakedefine01 MIOPEN_ENABLE_AI_KERNEL_TUNING
+#cmakedefine01 MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK
 
 // "_PACKAGE_" to avoid name contentions: the macros like
 // HIP_VERSION_MAJOR are defined in hip_version.h.
diff --git a/src/comgr.cpp b/src/comgr.cpp
index 3bf440462d..41af1d0e95 100644
--- a/src/comgr.cpp
+++ b/src/comgr.cpp
@@ -1003,15 +1003,9 @@ void BuildAsm(const std::string& name,
 
 #define WORKAROUND_ISSUE_HIPRTC_HIPRTC_HEADER_H 1 // See SWDEV-307838, issue #1648.
 #define WORKAROUND_ISSUE_1674 (HIP_PACKAGE_VERSION_FLAT >= 5003022305ULL)
-/// No assumption that HIP kernels are launched with uniform block size for backward compatibility
-/// SWDEV-413293 and https://reviews.llvm.org/D155213 effective HIP_FLAT_VERSION 500723302
-#ifndef _WIN32
-#define WORKAROUND_SWDEV_413293 (HIP_PACKAGE_VERSION_FLAT >= 5007023302ULL)
-#else
-/// Do not use on Windows. Compiler error message:
-///     '-fno-offload-uniform-block' - unknown command line option
-#define WORKAROUND_SWDEV_413293 0
-#endif
+
+// See WORKAROUND_SWDEV_413293 in ./CmakeLists.txt
+#define WORKAROUND_SWDEV_413293 MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK
 
 namespace hiprtc {