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 458c833)

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

@@ -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()
 

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.

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 {