Remove HIP PCH workaround

src/CMakeLists.txt

@@ -506,7 +506,6 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
         kernels/winograd/Conv_Winograd_Fury_v2_4_1_gfx11_1536vgprs_fp16_fp16acc_f2x3_c32_stride1.inc
         kernels/winograd/Conv_Winograd_Fury_v2_4_1_gfx11_1024vgprs_fp16_fp16acc_f2x3_c16_stride1.inc
         kernels/winograd/Conv_Winograd_Fury_v2_4_1_metadata.inc
-        kernels/workaround_issue_1431.hpp
         kernels/xform_bidirect_winograd_code.inc
         kernels/xform_data_filter.inc
         kernels/xform_kd_cov2.inc

src/comgr.cpp

@@ -101,23 +101,6 @@ MIOPEN_DECLARE_ENV_VAR_STR(ROCM_PATH)
 #error "AMD COMgr older than 1.7.0 is not supported"
 #endif
 
-#define COMGR_SUPPORTS_PCH (COMGR_VERSION >= 1008000)
-
-#if COMGR_SUPPORTS_PCH
-#if defined(__HIP_HAS_GET_PCH) && __HIP_HAS_GET_PCH
-#define HIP_SUPPORTS_PCH 1
-#else
-#define HIP_SUPPORTS_PCH 0
-#endif
-#endif // COMGR_SUPPORTS_PCH
-
-#define PCH_IS_SUPPORTED (COMGR_SUPPORTS_PCH && HIP_SUPPORTS_PCH)
-
-/// It seems like precompiled headers are built with "warpSize" fixed to 64.
-/// This leads to issues in HIP kernels that use "warpSize" on devices that
-/// have wavesize != 64 (currently gfx10 with default build settings).
-#define WORKAROUND_ISSUE_1431 PCH_IS_SUPPORTED
-
 #define EC_BASE(comgrcall, info, action)                                  \
     do                                                                    \
     {                                                                     \
@@ -589,14 +572,6 @@ static void SetIsaName(const ActionInfo& action,
     action.SetIsaName(isaName);
 }
 
-#if WORKAROUND_ISSUE_1431
-static inline bool IsWave64Enforced(const OptionList& opts)
-{
-    return std::any_of(
-        opts.begin(), opts.end(), [](const std::string& s) { return s == "-mwavefrontsize64"; });
-}
-#endif
-
 void BuildOcl(const std::string& name,
               std::string_view text,
               const std::string& options,
@@ -966,11 +941,6 @@ void BuildHip(const std::string& name,
 #if HIP_PACKAGE_VERSION_FLAT < 6001024000ULL && !defined(_WIN32)
         opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1");
 #endif
-#if WORKAROUND_ISSUE_1431
-        if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) &&
-           !miopen::comgr::IsWave64Enforced(opts))
-            opts.push_back("-DWORKAROUND_ISSUE_1431=1");
-#endif
 #if WORKAROUND_ISSUE_HIPRTC_HIPRTC_HEADER_H
         opts.push_back("-Wno-newline-eof");
         opts.push_back("-Wno-reserved-identifier");

src/composable_kernel/composable_kernel/include/tensor_operation/gridwise_generic_2d_reduction_direct_warpwise.hpp

@@ -26,9 +26,6 @@
 #ifndef CK_GRIDWISE_GENERIC_2D_REDUCTION_DIRECT_WARPWISE_HPP
 #define CK_GRIDWISE_GENERIC_2D_REDUCTION_DIRECT_WARPWISE_HPP
 
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "data_type.hpp"
 #include "reduction_common.hpp"
 #include "reduction_operator.hpp"

src/composable_kernel/composable_kernel/include/tensor_operation/reduction_functions_warpwise.hpp

@@ -26,9 +26,6 @@
 #ifndef CK_REDUCTION_FUNCTIONS_WARPWISE_HPP
 #define CK_REDUCTION_FUNCTIONS_WARPWISE_HPP
 
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "data_type.hpp"
 
 #include "reduction_common.hpp"

src/composable_kernel/composable_kernel/src/kernel_wrapper/gridwise_generic_reduction_first_call_warpwise_reduce_all_dims.cpp

@@ -23,9 +23,6 @@
  * SOFTWARE.
  *
  *******************************************************************************/
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "config.hpp"
 #include "number.hpp"
 #include "sequence.hpp"

src/composable_kernel/composable_kernel/src/kernel_wrapper/gridwise_generic_reduction_first_call_warpwise_reduce_partial_dims.cpp

@@ -23,9 +23,6 @@
  * SOFTWARE.
  *
  *******************************************************************************/
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "config.hpp"
 #include "number.hpp"
 #include "sequence.hpp"

src/composable_kernel/composable_kernel/src/kernel_wrapper/gridwise_generic_reduction_second_call_warpwise_reduce_all_dims.cpp

@@ -23,9 +23,6 @@
  * SOFTWARE.
  *
  *******************************************************************************/
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "config.hpp"
 #include "number.hpp"
 #include "sequence.hpp"

src/composable_kernel/composable_kernel/src/kernel_wrapper/gridwise_generic_reduction_second_call_warpwise_reduce_partial_dims.cpp

@@ -23,9 +23,6 @@
  * SOFTWARE.
  *
  *******************************************************************************/
-#ifdef WORKAROUND_ISSUE_1431
-#include "workaround_issue_1431.hpp"
-#endif
 #include "config.hpp"
 #include "number.hpp"
 #include "sequence.hpp"