[6.0.x] Fix for issue #2734: Detect if "-fno-offload-uniform-block" w…

…orks in HIP compiler. (#2741) * fix-issue-2734-rel-6.0 (01) Fix build with -DMIOPEN_USE_COMPOSABLEKERNEL=off for 6.0 release * fix-issue-2734-rel-6.0 (02) Use "-fno-offload-uniform-block" only if HIP compiler supports it. Resolves #2734.
ROCm · Feb 14, 2024 · 9110d44 · 9110d44
1 parent a55aaa0
commit 9110d44
Show file tree

Hide file tree

Showing 5 changed files with 19 additions and 4 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -252,7 +252,17 @@ set_var_to_condition(MIOPEN_USE_HIPRTC_DEFAULT ${MIOPEN_USE_COMGR} AND (${MIOPEN
 option(MIOPEN_USE_HIPRTC "Use HIPRTC to build HIP kernels instead of COMGR" ${MIOPEN_USE_HIPRTC_DEFAULT})
 
 # WORKAROUND_SWDEV_413293
-if(${MIOPEN_hip_VERSION_FLAT} GREATER_EQUAL 500723302)
+if(HAS_HIP)
+    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
@@ -51,6 +51,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
@@ -1323,7 +1323,7 @@ void BuildHip(const std::string& name,
         opts.push_back("-Wno-cuda-compat");
         opts.push_back("-fno-gpu-rdc");
         opts.push_back("-O3");
-#if WORKAROUND_SWDEV_413293
+#if WORKAROUND_SWDEV_413293 && MIOPEN_HIP_COMPILER_HAS_OPTION_OFFLOAD_UNIFORM_BLOCK
         opts.push_back("-fno-offload-uniform-block");
 #endif
         if(std::none_of(opts.begin(), opts.end(), [](const std::string& s) {

diff --git a/src/solver/batchnorm/backward_ck.cpp b/src/solver/batchnorm/backward_ck.cpp
@@ -161,7 +161,7 @@ bool BnCKBwdBackward::IsApplicable(const ExecutionContext& context,
 {
 #if !MIOPEN_BACKEND_HIP || !MIOPEN_USE_COMPOSABLEKERNEL
     std::ignore = context;
-    std::ignore = fdesc_problem;
+    std::ignore = bn_problem;
     return false;
 #else
     if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_CK_BN_BACK{}))
@@ -191,6 +191,7 @@ bool BnCKBwdBackward::IsApplicable(const ExecutionContext& context,
 #endif
 }
 
+#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
 template <typename XDataType,
           typename DxDataType,
           typename DyDataType,
@@ -220,6 +221,7 @@ ConvSolution MakeAnyInvokerFactory(const miopen::batchnorm::ProblemDescription&
                                  CKArgsBNormBwd,
                                  miopen::batchnorm::BwdInvokeParams>(bn_problem, kernel_id);
 }
+#endif
 
 ConvSolution BnCKBwdBackward::GetSolution(
     [[maybe_unused]] const ExecutionContext& context,

diff --git a/src/solver/batchnorm/forward_training_ck.cpp b/src/solver/batchnorm/forward_training_ck.cpp
@@ -156,7 +156,7 @@ bool BnCKFwdTraining::IsApplicable(const ExecutionContext& context,
 {
 #if !MIOPEN_BACKEND_HIP || !MIOPEN_USE_COMPOSABLEKERNEL
     std::ignore = context;
-    std::ignore = fdesc_problem;
+    std::ignore = bn_problem;
     return false;
 #else
     if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_CK_BN_FWD_TRAINING{}))
@@ -183,6 +183,7 @@ bool BnCKFwdTraining::IsApplicable(const ExecutionContext& context,
 #endif
 }
 
+#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
 template <typename XDataType,
           typename YDataType,
           typename AccDataType,
@@ -209,6 +210,7 @@ ConvSolution MakeAnyInvokerFactory(const miopen::batchnorm::ProblemDescription&
                                  CKArgsBNormFwdTraining,
                                  miopen::batchnorm::InvokeParams>(bn_problem, kernel_id);
 }
+#endif
 
 ConvSolution BnCKFwdTraining::GetSolution(
     [[maybe_unused]] const ExecutionContext& context,