[ROCm 6.1][hipRTC] Fix build failures. [quality] Reorg standard inclu…

…des in HIP sources. (#2637)

CMakeLists.txt

@@ -383,7 +383,7 @@ if(MIOPEN_USE_HIPRTC)
         message(FATAL_ERROR "HIPRTC can be used only together with COMGR")
     endif()
     find_package(hiprtc REQUIRED)
-    message(STATUS "Build with HIPRTC")
+    message(STATUS "Build with HIPRTC ${hiprtc_VERSION}")
 endif()
 
 option(Boost_USE_STATIC_LIBS "Use boost static libraries" ON)

driver/random.hpp

@@ -109,7 +109,9 @@ inline T gen_subnorm()
     if constexpr(!std::is_integral_v<T> && !std::is_same_v<T, double> &&
                  details::has_digits<T>::value)
     {
-        using BitType = std::conditional_t<sizeof(T) == 2, uint16_t, uint32_t>;
+        using BitType = std::conditional_t<sizeof(T) == 1,
+                                           uint8_t,
+                                           std::conditional_t<sizeof(T) == 2, uint16_t, uint32_t>>;
         static_assert(sizeof(T) == sizeof(BitType));
 
         // -1 because ::digits counts the first implicit digit

src/CMakeLists.txt

@@ -381,23 +381,27 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
         kernels/Conv_Winograd_v30_3_1_gfx11_fp32_f3x2_stride1.inc
         kernels/Conv_Winograd_v30_3_1_gfx11_fp32_f3x2_stride2.inc
         kernels/Conv_Winograd_v30_3_1_metadata.inc
-        kernels/xform_bidirect_winograd_code.inc
-        kernels/rocm_version.inc
-        kernels/inst_wrappers.inc
+        kernels/bfloat16_dev.hpp
         kernels/conv_common.inc
-        kernels/utilities.inc
-        kernels/xform_data_filter.inc
-        kernels/xform_kd_cov2.inc
-        kernels/xform_metadata.inc
-        kernels/neuron.inc
         kernels/conv_sizes.inc
-        kernels/gpr_alloc.inc
-        kernels/bfloat16_dev.hpp
         kernels/float_types.h
-        kernels/workaround_issue_1431.hpp
+        kernels/gpr_alloc.inc
         kernels/hip_f8_impl.hpp
         kernels/hip_float8.hpp
+        kernels/inst_wrappers.inc
+        kernels/miopen_cstdint.hpp
+        kernels/miopen_limits.hpp
+        kernels/miopen_type_traits.hpp
+        kernels/miopen_utility.hpp
+        kernels/neuron.inc
+        kernels/rocm_version.inc
         kernels/stride_array.hpp
+        kernels/utilities.inc
+        kernels/workaround_issue_1431.hpp
+        kernels/xform_bidirect_winograd_code.inc
+        kernels/xform_data_filter.inc
+        kernels/xform_kd_cov2.inc
+        kernels/xform_metadata.inc
         )
 
     set(MIOPEN_KERNELS

src/comgr.cpp

@@ -1292,8 +1292,7 @@ void BuildHip(const std::string& name,
         auto opts =
             miopen::SplitSpaceSeparated(options, miopen::comgr::compiler::lc::GetOptionsNoSplit());
         compiler::lc::RemoveOptionsUnwanted(opts);
-        opts.push_back("-DWORKAROUND_ISSUE_HIPRTC_TRUE_TYPE"); // Workaround for SWDEV-308073
-#if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL
+#if HIP_PACKAGE_VERSION_MAJOR < 6
         opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround?
 #endif
         opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround?
@@ -1302,7 +1301,11 @@ void BuildHip(const std::string& name,
             opts.push_back("-DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1");
 #endif
         opts.push_back("-DHIP_PACKAGE_VERSION_FLAT=" + std::to_string(HIP_PACKAGE_VERSION_FLAT));
-        opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS=1");
+        opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS");
+        /// For now, use only standard <limits> to avoid possibility of
+        /// correctnes or performance regressions.
+        /// \todo Test and enable "custom" local implementation.
+        opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1");
 #if WORKAROUND_ISSUE_1431
         if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) &&
            !miopen::comgr::IsWave64Enforced(opts))

src/composable_kernel/composable_kernel/include/utility/array.hpp

@@ -4,34 +4,7 @@
 #include "functional2.hpp"
 #include "sequence.hpp"
 
-#ifdef __HIPCC_RTC__
-#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
-/// We need <utility> for std::forward. In some cases, it includes <type_traits>
-/// (this is against the Standard, but it doesn't matter in this case).
-/// But <type_traits> also defines std::true_type, per Standard.
-/// However the latter definition conflicts with
-/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h,
-/// which defines std::true_type as well (which is wrong).
-
-namespace std {
-
-template <typename T>
-constexpr T&& forward(typename remove_reference<T>::type& t_) noexcept
-{
-    return static_cast<T&&>(t_);
-}
-
-template <typename T>
-constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
-{
-    return static_cast<T&&>(t_);
-}
-
-} // namespace std
-#else
-#include <utility> // std::forward
-#endif
-#endif // __HIPCC_RTC__
+#include "miopen_utility.hpp" // std::forward
 
 namespace ck {
 

src/composable_kernel/composable_kernel/include/utility/config.hpp

@@ -6,6 +6,7 @@
 #include "hip/hip_fp16.h"
 #endif
 #include "bfloat16_dev.hpp"
+#include "miopen_cstdint.hpp"
 
 // "Constant" address space for kernel parameter
 #define CONSTANT __attribute__((address_space(4)))

src/composable_kernel/composable_kernel/include/utility/data_type.hpp

@@ -3,21 +3,8 @@
 
 #include "statically_indexed_array.hpp"
 
-#ifdef __HIPCC_RTC__
-#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
-/// Definitions from <cstdint>, <cmath> conflict with
-/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h.
-
-typedef signed char int8_t;
-typedef signed short int16_t;
-typedef float float_t;
-#include <limits> // std::numeric_limits
-
-#else
-#include <cstdint> // int8_t, int16_t
-#include <cmath>   // float_t
-#endif
-#endif // __HIPCC_RTC__
+#include "miopen_cstdint.hpp"
+#include "miopen_limits.hpp"
 
 namespace ck {
 
@@ -978,7 +965,7 @@ struct inner_product_with_conversion
         return acc;
     }
 
-    __device__ T operator()(float_t a, float_t b) const { return convert(a) * convert(b); }
+    __device__ T operator()(float a, float b) const { return convert(a) * convert(b); }
 
     __device__ T operator()(int8x4_t a, int8x4_t b) const
     {

src/composable_kernel/composable_kernel/include/utility/enable_if.hpp

@@ -1,6 +1,8 @@
 #ifndef CK_ENABLE_IF_HPP
 #define CK_ENABLE_IF_HPP
 
+#include "miopen_type_traits.hpp"
+
 namespace ck {
 
 template <bool B, typename T = void>

src/composable_kernel/composable_kernel/include/utility/magic_division.hpp

@@ -6,6 +6,7 @@
 #include "number.hpp"
 #include "type.hpp"
 #include "tuple.hpp"
+#include "miopen_cstdint.hpp"
 
 namespace ck {
 

src/composable_kernel/composable_kernel/include/utility/type.hpp

@@ -3,84 +3,7 @@
 
 #include "integral_constant.hpp"
 #include "enable_if.hpp"
-
-#ifdef __HIPCC_RTC__
-#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
-/// We need <type_traits> for std::remove_reference and std::remove_cv.
-/// But <type_traits> also defines std::true_type, per Standard.
-/// However the latter definition conflicts with
-/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h,
-/// which defines std::true_type as well (which is wrong).
-
-namespace std {
-
-template <class T>
-struct remove_reference
-{
-    typedef T type;
-};
-template <class T>
-struct remove_reference<T&>
-{
-    typedef T type;
-};
-template <class T>
-struct remove_reference<T&&>
-{
-    typedef T type;
-};
-
-template <class T>
-using remove_reference_t = typename remove_reference<T>::type;
-
-template <class T>
-struct remove_const
-{
-    typedef T type;
-};
-template <class T>
-struct remove_const<const T>
-{
-    typedef T type;
-};
-
-template <class T>
-struct remove_volatile
-{
-    typedef T type;
-};
-template <class T>
-struct remove_volatile<volatile T>
-{
-    typedef T type;
-};
-
-template <class T>
-struct remove_cv
-{
-    typedef typename remove_volatile<typename remove_const<T>::type>::type type;
-};
-
-template <class T>
-struct is_pointer_helper : std::false_type
-{
-};
-
-template <class T>
-struct is_pointer_helper<T*> : std::true_type
-{
-};
-
-template <class T>
-struct is_pointer : is_pointer_helper<typename std::remove_cv<T>::type>
-{
-};
-
-} // namespace std
-#else
-#include <type_traits> // std::remove_reference, std::remove_cv, is_pointer
-#endif
-#endif // __HIPCC_RTC__
+#include "miopen_type_traits.hpp"
 
 namespace ck {
 

src/gemm_v2.cpp

@@ -119,8 +119,6 @@ static inline rocblas_computetype rocBlasComputeType_ex3(const miopen::GemmDescr
 
 static inline rocblas_datatype rocBlasComputeType(const miopen::GemmDescriptor& desc)
 {
-    // Complex compute types are only supported in newer version of the API
-    assert(desc.dataType == desc.a_cast_type && desc.dataType == desc.b_cast_type);
     if(desc.dataType == miopenInt8)
         return rocblas_datatype::rocblas_datatype_i32_r;
     else

src/kernels/MIOpenCheckNumerics.cpp

@@ -30,51 +30,11 @@
 
 #include <hip/hip_bfloat16.h>
 
-// Copied over from naive_conv.cpp
-#ifdef __HIPCC_RTC__
-#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
-/// Definitions from <cstdint>, <cmath> conflict with
-/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h.
-
-typedef unsigned char uint8_t;
-typedef signed char int8_t;
-typedef signed short int16_t;
-typedef unsigned short uint16_t;
-typedef float float_t;
-
-// std::conditional requires type_traits which has a few other things
-// which result in collision with amd_hip_vector_types.h
-
-namespace std {
-template <bool predicate, typename X, typename Y>
-struct conditional;
-
-template <typename X, typename Y>
-struct conditional<true, X, Y>
-{
-    using type = X;
-};
-
-template <typename X, typename Y>
-struct conditional<false, X, Y>
-{
-    using type = Y;
-};
-
-template <bool predicate, typename X, typename Y>
-using conditional_t = typename conditional<predicate, X, Y>::type;
-} // namespace std
-#else
-#include <cstdint> // int8_t, int16_t
-#include <cmath>   // float_t
-#endif
-#endif // __HIPCC_RTC__
-
-#include <limits> // std::numeric_limits
-
 #define MIOPEN_ENABLE_F8_DEVICE_CODE 1
 #include "hip_float8.hpp"
 
+#include "miopen_limits.hpp"
+
 struct Numerics
 {
     float sum;

src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp

@@ -28,6 +28,8 @@
 #include <hip/hip_fp16.h>
 #endif
 
+#include "miopen_cstdint.hpp"
+
 #ifndef BATCHED_TRANSPOSE_OCCUPANCY
 #define BATCHED_TRANSPOSE_OCCUPANCY 4
 #endif

src/kernels/gpu_general_tensor_reorder_kernel/general_tensor_reorder_kernel_util.hpp

@@ -24,34 +24,23 @@
  *
  *******************************************************************************/
 #ifndef GENERAL_TENSOR_REORDER_UTIL_HPP
-#ifdef __HIPCC_RTC__
-#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
-/// Definitions from <cstdint>, <cmath> conflict with
-/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h.
-
-typedef signed char int8_t;
-typedef signed short int16_t;
-typedef unsigned int uint32_t;
-
-#else
-#include <cstdint> // int8_t, int16_t
-#endif
-#endif // __HIPCC_RTC__
+#define GENERAL_TENSOR_REORDER_UTIL_HPP
 
 #ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS
 #include <hip/hip_runtime.h>
 #include <hip/hip_fp16.h>
 #endif
 
+#include "miopen_cstdint.hpp"
+
 #ifndef TENSOR_REORDER_OCCUPANCY
 #define TENSOR_REORDER_OCCUPANCY 4
 #endif
-#define GENERAL_TENSOR_REORDER_UTIL_HPP
 
 template <int... Is>
 struct order
 {
-    static constexpr std::size_t m_size = sizeof...(Is);
+    static constexpr size_t m_size = sizeof...(Is);
     // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
     static constexpr int m_data[m_size + 1] = {Is..., 0};