diff --git a/CMakeLists.txt b/CMakeLists.txt index aebd984989..4f38e3671a 100644 --- 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 index 74b0185fde..3d1d44faf0 100644 --- 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 index 18f41c862d..e0f07516a4 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -1318,7 +1318,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) {