SWDEV-419480 - HIPRTC implementation of standard headers shouldn't be in std namespace

HIPRTC implementation of few type traits are defined in std
namespace which causes conflicts with the std C++ headers when
apps try to include those std headers. Instead, use the
specific implementation of these headers from __hip_internal
namespace to avoid this.

Change-Id: Ia358df159af8783d1eba9a19cd458e192bf7e9e6


[ROCm/clr commit: 353dbe6e3b]

projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h

@@ -84,13 +84,13 @@ THE SOFTWARE.
         #include "amd_device_functions.h"
         #include "amd_warp_functions.h"
       #endif
-        namespace std
+        namespace __hip_internal
         {
-            template<> struct is_floating_point<_Float16> : std::true_type {};
+            template<> struct is_floating_point<_Float16> : __hip_internal::true_type {};
         }
 
         template<bool cond, typename T = void>
-        using Enable_if_t = typename std::enable_if<cond, T>::type;
+        using Enable_if_t = typename __hip_internal::enable_if<cond, T>::type;
 
         // BEGIN STRUCT __HALF
         struct __half {
@@ -112,7 +112,7 @@ THE SOFTWARE.
                 __half(decltype(data) x) : data{x} {}
                 template<
                     typename T,
-                    Enable_if_t<std::is_floating_point<T>{}>* = nullptr>
+                    Enable_if_t<__hip_internal::is_floating_point<T>{}>* = nullptr>
                 __HOST_DEVICE__
                 __half(T x) : data{static_cast<_Float16>(x)} {}
             #endif
@@ -126,7 +126,7 @@ THE SOFTWARE.
             // CREATORS - DEVICE ONLY
             #if !defined(__HIP_NO_HALF_CONVERSIONS__)
                 template<
-                    typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
+                    typename T, Enable_if_t<__hip_internal::is_integral<T>{}>* = nullptr>
                 __HOST_DEVICE__
                 __half(T x) : data{static_cast<_Float16>(x)} {}
             #endif
@@ -171,7 +171,7 @@ THE SOFTWARE.
             #if !defined(__HIP_NO_HALF_CONVERSIONS__)
                 template<
                     typename T,
-                    Enable_if_t<std::is_floating_point<T>{}>* = nullptr>
+                    Enable_if_t<__hip_internal::is_floating_point<T>{}>* = nullptr>
                 __HOST_DEVICE__
                 __half& operator=(T x)
                 {
@@ -183,7 +183,7 @@ THE SOFTWARE.
             // MANIPULATORS - DEVICE ONLY
             #if !defined(__HIP_NO_HALF_CONVERSIONS__)
                 template<
-                    typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
+                    typename T, Enable_if_t<__hip_internal::is_integral<T>{}>* = nullptr>
                 __device__
                 __half& operator=(T x)
                 {
@@ -241,7 +241,7 @@ THE SOFTWARE.
             #if !defined(__HIP_NO_HALF_CONVERSIONS__)
                 template<
                     typename T,
-                    Enable_if_t<std::is_floating_point<T>{}>* = nullptr>
+                    Enable_if_t<__hip_internal::is_floating_point<T>{}>* = nullptr>
                 __HOST_DEVICE__
                 operator T() const { return data; }
             #endif
@@ -255,7 +255,7 @@ THE SOFTWARE.
 
             #if !defined(__HIP_NO_HALF_CONVERSIONS__)
                 template<
-                    typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
+                    typename T, Enable_if_t<__hip_internal::is_integral<T>{}>* = nullptr>
                 __HOST_DEVICE__
                 operator T() const { return data; }
             #endif

projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h

@@ -90,14 +90,10 @@ size_t amd_dbgapi_get_build_id();
 #else
 #if !__HIP_NO_STD_DEFS__
 typedef unsigned int uint32_t;
-typedef unsigned long long uint64_t;
 typedef signed int int32_t;
-typedef signed long long int64_t;
 namespace std {
 using ::uint32_t;
-using ::uint64_t;
 using ::int32_t;
-using ::int64_t;
 }
 #endif // __HIP_NO_STD_DEFS__
 #endif // !defined(__HIPCC_RTC__)

projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -48,88 +48,6 @@ THE SOFTWARE.
     #include <array>
     #include <iosfwd>
     #include <type_traits>
-#else
-namespace std {
-using ::size_t;
-
-template <class _Tp, _Tp __v> struct integral_constant {
-  static constexpr const _Tp value = __v;
-  typedef _Tp value_type;
-  typedef integral_constant type;
-  constexpr operator value_type() const { return value; }
-  constexpr value_type operator()() const { return value; }
-};
-template <class _Tp, _Tp __v> constexpr const _Tp integral_constant<_Tp, __v>::value;
-
-typedef integral_constant<bool, true> true_type;
-typedef integral_constant<bool, false> false_type;
-
-template <bool B> using bool_constant = integral_constant<bool, B>;
-typedef bool_constant<true> true_type;
-typedef bool_constant<false> false_type;
-
-template <bool __B, class __T = void> struct enable_if {};
-template <class __T> struct enable_if<true, __T> { typedef __T type; };
-
-template<bool _B> struct true_or_false_type : public false_type {};
-template<> struct true_or_false_type<true> : public true_type {};
-
-template <class _Tp> struct is_integral : public false_type {};
-template <> struct is_integral<bool> : public true_type {};
-template <> struct is_integral<char> : public true_type {};
-template <> struct is_integral<signed char> : public true_type {};
-template <> struct is_integral<unsigned char> : public true_type {};
-template <> struct is_integral<wchar_t> : public true_type {};
-template <> struct is_integral<short> : public true_type {};
-template <> struct is_integral<unsigned short> : public true_type {};
-template <> struct is_integral<int> : public true_type {};
-template <> struct is_integral<unsigned int> : public true_type {};
-template <> struct is_integral<long> : public true_type {};
-template <> struct is_integral<unsigned long> : public true_type {};
-template <> struct is_integral<long long> : public true_type {};
-template <> struct is_integral<unsigned long long> : public true_type {};
-
-template <class _Tp> struct is_arithmetic : public false_type {};
-template <> struct is_arithmetic<bool> : public true_type {};
-template <> struct is_arithmetic<char> : public true_type {};
-template <> struct is_arithmetic<signed char> : public true_type {};
-template <> struct is_arithmetic<unsigned char> : public true_type {};
-template <> struct is_arithmetic<wchar_t> : public true_type {};
-template <> struct is_arithmetic<short> : public true_type {};
-template <> struct is_arithmetic<unsigned short> : public true_type {};
-template <> struct is_arithmetic<int> : public true_type {};
-template <> struct is_arithmetic<unsigned int> : public true_type {};
-template <> struct is_arithmetic<long> : public true_type {};
-template <> struct is_arithmetic<unsigned long> : public true_type {};
-template <> struct is_arithmetic<long long> : public true_type {};
-template <> struct is_arithmetic<unsigned long long> : public true_type {};
-template <> struct is_arithmetic<float> : public true_type {};
-template <> struct is_arithmetic<double> : public true_type {};
-
-template<typename _Tp> struct is_floating_point : public false_type {};
-template<> struct is_floating_point<float> : public true_type {};
-template<> struct is_floating_point<double> : public true_type {};
-template<> struct is_floating_point<long double> : public true_type {};
-
-template <typename __T, typename __U> struct is_same : public false_type {};
-template <typename __T> struct is_same<__T, __T> : public true_type {};
-
-template<typename _Tp, bool = is_arithmetic<_Tp>::value>
-  struct is_signed : public false_type {};
-template<typename _Tp>
-  struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {};
-
-template <class _T1, class _T2> struct is_convertible
-  : public true_or_false_type<__is_convertible_to(_T1, _T2)> {};
-
-template<typename _CharT> struct char_traits;
-template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_istream;
-template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_ostream;
-typedef basic_istream<char> istream;
-typedef basic_ostream<char> ostream;
-
-template <typename __T> struct is_scalar : public integral_constant<bool, __is_scalar(__T)> {};
-} // Namespace std.
 #endif // defined(__HIPCC_RTC__)
 
     namespace hip_impl {
@@ -275,7 +193,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
             template<
                 typename U = T,
-                typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_signed<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_ operator-() const noexcept
             {
@@ -286,7 +204,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_ operator~() const noexcept
             {
@@ -296,7 +214,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator%=(const Native_vec_& x_) noexcept
             {
@@ -305,7 +223,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator^=(const Native_vec_& x_) noexcept
             {
@@ -314,7 +232,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator|=(const Native_vec_& x_) noexcept
             {
@@ -323,7 +241,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator&=(const Native_vec_& x_) noexcept
             {
@@ -332,7 +250,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator>>=(const Native_vec_& x_) noexcept
             {
@@ -341,7 +259,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
             }
             template<
                 typename U = T,
-                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+                typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
             __HOST_DEVICE__
             Native_vec_& operator<<=(const Native_vec_& x_) noexcept
             {
@@ -448,8 +366,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         HIP_vector_type() = default;
         template<
             typename U,
-            typename std::enable_if<
-                std::is_convertible<U, T>::value>::type* = nullptr>
+            typename __hip_internal::enable_if<
+                __hip_internal::is_convertible<U, T>::value>::type* = nullptr>
         __HOST_DEVICE__
         explicit
         constexpr
@@ -458,7 +376,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         {}
         template< // TODO: constrain based on type as well.
             typename... Us,
-            typename std::enable_if<
+            typename __hip_internal::enable_if<
                 (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
         __HOST_DEVICE__
         constexpr
@@ -514,8 +432,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         }
         template<
             typename U,
-            typename std::enable_if<
-                std::is_convertible<U, T>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<
+                __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator+=(U x) noexcept
         {
@@ -530,8 +448,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         }
         template<
             typename U,
-            typename std::enable_if<
-                std::is_convertible<U, T>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<
+                __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator-=(U x) noexcept
         {
@@ -553,8 +471,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U,
-            typename std::enable_if<
-                std::is_convertible<U, T>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<
+                __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator*=(U x) noexcept
         {
@@ -575,8 +493,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         }
         template<
             typename U,
-            typename std::enable_if<
-                std::is_convertible<U, T>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<
+                __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator/=(U x) noexcept
         {
@@ -585,7 +503,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_signed<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type operator-() const noexcept
         {
@@ -596,7 +514,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type operator~() const noexcept
         {
@@ -607,7 +525,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
         {
@@ -617,7 +535,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
         {
@@ -627,7 +545,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
         {
@@ -637,7 +555,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
         {
@@ -647,7 +565,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
         {
@@ -657,7 +575,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 
         template<
             typename U = T,
-            typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            typename __hip_internal::enable_if<__hip_internal::is_integral<U>{}>::type* = nullptr>
         __HOST_DEVICE__
         HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
         {
@@ -825,7 +743,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -838,7 +756,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -851,7 +769,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -864,7 +782,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -877,7 +795,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -890,7 +808,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -903,7 +821,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -916,7 +834,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -929,7 +847,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -942,7 +860,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -955,7 +873,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -968,7 +886,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -981,7 +899,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -994,7 +912,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -1007,7 +925,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -1020,7 +938,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     template<
         typename T,
         unsigned int n,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -1033,7 +951,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -1046,8 +964,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         typename T,
         unsigned int n,
         typename U,
-        typename std::enable_if<std::is_arithmetic<U>::value>::type,
-        typename std::enable_if<std::is_integral<T>{}>* = nullptr>
+        typename __hip_internal::enable_if<__hip_internal::is_arithmetic<U>::value>::type,
+        typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
     __HOST_DEVICE__
     inline
     constexpr
@@ -1061,28 +979,28 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
      * Map HIP_vector_type<U, rankU> to HIP_vector_type<T, rankT>
      */
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 1 && rankU >= 1),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 1 && rankU >= 1),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT>(static_cast<T>(u.x));
     };
 
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU == 1),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU == 1),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(0));
     };
 
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU >= 2),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU >= 2),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(u.y));
     };
 
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 1),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 1),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(0),
@@ -1090,7 +1008,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     };
 
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 2),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 2),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT>(static_cast<T>(u.x), static_cast<T>(u.y),
@@ -1098,7 +1016,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
     };
 
     template <typename T, unsigned int rankT, typename U, unsigned int rankU>
-    __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 4),
+    __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 4),
                                                             const HIP_vector_type<T, rankT>>::type
     __hipMapVector(const HIP_vector_type<U, rankU>& u) {
       return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(u.y),

projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h

@@ -111,7 +111,7 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
         int boundaryMode = hipBoundaryModeZero) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
@@ -122,7 +122,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -132,7 +132,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -142,7 +142,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -152,7 +152,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
@@ -162,7 +162,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
@@ -172,7 +172,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -182,7 +182,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -192,7 +192,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -202,7 +202,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -212,7 +212,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -222,7 +222,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
     x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -232,7 +232,7 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
         int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT
@@ -243,7 +243,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
         int layer) {
     __HIP_SURFACE_OBJECT_PARAMETERS_INIT

projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h

@@ -70,18 +70,20 @@ namespace cooperative_groups {
 
 /* Global scope */
 template <unsigned int size>
-using is_power_of_2 = std::integral_constant<bool, (size & (size - 1)) == 0>;
+using is_power_of_2 = __hip_internal::integral_constant<bool, (size & (size - 1)) == 0>;
 
 template <unsigned int size>
-using is_valid_wavefront = std::integral_constant<bool, (size <= __AMDGCN_WAVEFRONT_SIZE)>;
+using is_valid_wavefront = __hip_internal::integral_constant<bool, (size <= __AMDGCN_WAVEFRONT_SIZE)>;
 
 template <unsigned int size>
 using is_valid_tile_size =
-    std::integral_constant<bool, is_power_of_2<size>::value && is_valid_wavefront<size>::value>;
+    __hip_internal::integral_constant<bool, is_power_of_2<size>::value &&
+          is_valid_wavefront<size>::value>;
 
 template <typename T>
 using is_valid_type =
-    std::integral_constant<bool, std::is_integral<T>::value || std::is_floating_point<T>::value>;
+    __hip_internal::integral_constant<bool, __hip_internal::is_integral<T>::value ||
+          __hip_internal::is_floating_point<T>::value>;
 
 namespace internal {
 

projects/clr/hipamd/include/hip/amd_detail/host_defines.h

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -33,8 +33,7 @@ THE SOFTWARE.
 #define GENERIC_GRID_LAUNCH 1
 #endif
 
-#if defined(__clang__) && defined(__HIP__)
-
+#if defined(__cplusplus)
 namespace __hip_internal {
 typedef unsigned char uint8_t;
 typedef unsigned short uint16_t;
@@ -112,6 +111,57 @@ template<typename _Tp, bool = is_arithmetic<_Tp>::value>
 template<typename _Tp>
   struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {};
 
+template<class T>
+    auto test_returnable(int) -> decltype(
+        void(static_cast<T(*)()>(nullptr)), true_type{});
+template<class>
+    auto test_returnable(...) -> false_type;
+
+template<class T>
+    struct type_identity { using type = T; };
+
+template<class T> // Note that `cv void&` is a substitution failure
+    auto try_add_lvalue_reference(int) -> type_identity<T&>;
+template<class T> // Handle T = cv void case
+    auto try_add_lvalue_reference(...) -> type_identity<T>;
+
+template<class T>
+    auto try_add_rvalue_reference(int) -> type_identity<T&&>;
+template<class T>
+    auto try_add_rvalue_reference(...) -> type_identity<T>;
+
+template<class T>
+struct add_lvalue_reference
+    : decltype(try_add_lvalue_reference<T>(0)) {};
+
+template<class T>
+struct add_rvalue_reference
+    : decltype(try_add_rvalue_reference<T>(0)) {};
+
+template<typename T>
+typename add_rvalue_reference<T>::type declval() noexcept;
+
+template<class From, class To>
+    auto test_implicitly_convertible(int) -> decltype(
+        void(declval<void(&)(To)>()(declval<From>())), true_type{});
+
+template<class, class>
+    auto test_implicitly_convertible(...) -> false_type;
+
+template<class T> struct remove_cv { typedef T type; };
+template<class T> struct remove_cv<const T> { typedef T type; };
+template<class T> struct remove_cv<volatile T> { typedef T type; };
+template<class T> struct remove_cv<const volatile T> { typedef T type; };
+
+template<class T>
+struct is_void : public is_same<void, typename remove_cv<T>::type> {};
+
+template<class From, class To>
+struct is_convertible : public integral_constant<bool,
+    (decltype(test_returnable<To>(0))::value &&
+     decltype(test_implicitly_convertible<From, To>(0))::value) ||
+    (is_void<From>::value && is_void<To>::value)> {};
+
 template<typename _CharT> struct char_traits;
 template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_istream;
 template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_ostream;
@@ -136,7 +186,9 @@ typedef __hip_internal::int8_t __hip_int8_t;
 typedef __hip_internal::int16_t __hip_int16_t;
 typedef __hip_internal::int32_t __hip_int32_t;
 typedef __hip_internal::int64_t __hip_int64_t;
+#endif // defined(__cplusplus)
 
+#if defined(__clang__) && defined(__HIP__)
 #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 #define __host__ __attribute__((host))
 #define __device__ __attribute__((device))
@@ -175,6 +227,6 @@ typedef __hip_internal::int64_t __hip_int64_t;
 #define __constant__
 
 #define __hip_img_chk__
-#endif
+#endif // defined(__clang__) && defined(__HIP__)
 
 #endif

projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h

@@ -51,13 +51,13 @@ template<typename T>
 struct __hip_is_tex_surf_scalar_channel_type
 {
     static constexpr bool value =
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value ||
-        std::is_same<T, int>::value ||
-        std::is_same<T, unsigned int>::value ||
-        std::is_same<T, float>::value;
+        __hip_internal::is_same<T, char>::value ||
+        __hip_internal::is_same<T, unsigned char>::value ||
+        __hip_internal::is_same<T, short>::value ||
+        __hip_internal::is_same<T, unsigned short>::value ||
+        __hip_internal::is_same<T, int>::value ||
+        __hip_internal::is_same<T, unsigned int>::value ||
+        __hip_internal::is_same<T, float>::value;
 };
 
 template<typename T>
@@ -83,10 +83,10 @@ template<typename T>
 struct __hip_is_tex_normalized_channel_type
 {
     static constexpr bool value =
-        std::is_same<T, char>::value ||
-        std::is_same<T, unsigned char>::value ||
-        std::is_same<T, short>::value ||
-        std::is_same<T, unsigned short>::value;
+        __hip_internal::is_same<T, char>::value ||
+        __hip_internal::is_same<T, unsigned char>::value ||
+        __hip_internal::is_same<T, short>::value ||
+        __hip_internal::is_same<T, unsigned short>::value;
 };
 
 template<
@@ -107,7 +107,7 @@ template <
     typename Enable = void>
 struct __hip_tex_ret
 {
-    static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
+    static_assert(__hip_internal::is_same<Enable, void>::value, "Invalid channel type!");
 };
 
 /*
@@ -115,7 +115,7 @@ struct __hip_tex_ret
  */
 template<typename T, typename U>
 __forceinline__ __device__
-typename std::enable_if<
+typename __hip_internal::enable_if<
   __hip_is_tex_surf_scalar_channel_type<T>::value, const T>::type
 __hipMapFrom(const U &u) {
   if constexpr (sizeof(T) < sizeof(float)) {
@@ -138,7 +138,7 @@ __hipMapFrom(const U &u) {
  */
 template<typename T, typename U>
 __forceinline__ __device__
-typename std::enable_if<
+typename __hip_internal::enable_if<
   __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const T>::type
 __hipMapFrom(const U &u) {
   if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
@@ -161,7 +161,7 @@ __hipMapFrom(const U &u) {
  */
 template<typename U, typename T>
 __forceinline__ __device__
-typename std::enable_if<
+typename __hip_internal::enable_if<
 __hip_is_tex_surf_scalar_channel_type<T>::value, const U>::type
 __hipMapTo(const T &t) {
   if constexpr (sizeof(T) < sizeof(float)) {
@@ -186,7 +186,7 @@ __hipMapTo(const T &t) {
  */
 template<typename U, typename T>
 __forceinline__ __device__
-typename std::enable_if<
+typename __hip_internal::enable_if<
   __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const U>::type
 __hipMapTo(const T &t) {
   if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
@@ -215,7 +215,7 @@ template <typename T>
 struct __hip_tex_ret<
     T,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
 {
     using type = T;
 };
@@ -226,7 +226,7 @@ template<
 struct __hip_tex_ret<
     HIP_vector_type<T, rank>,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
 };
@@ -235,7 +235,7 @@ template<typename T>
 struct __hip_tex_ret<
     T,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
 {
     using type = float;
 };
@@ -246,7 +246,7 @@ template<
 struct __hip_tex_ret<
     HIP_vector_type<T, rank>,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
 };
@@ -438,7 +438,7 @@ template <
     typename Enable = void>
 struct __hip_tex2dgather_ret
 {
-    static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
+    static_assert(__hip_internal::is_same<Enable, void>::value, "Invalid channel type!");
 };
 
 template <
@@ -450,7 +450,7 @@ template <typename T>
 struct __hip_tex2dgather_ret<
     T,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
 {
     using type = HIP_vector_type<T, 4>;
 };
@@ -461,7 +461,7 @@ template<
 struct __hip_tex2dgather_ret<
     HIP_vector_type<T, rank>,
     hipReadModeElementType,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
 {
     using type = HIP_vector_type<T, 4>;
 };
@@ -470,7 +470,7 @@ template <typename T>
 struct __hip_tex2dgather_ret<
     T,
     hipReadModeNormalizedFloat,
-    typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
+    typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
 {
     using type = float4;
 };

projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h

@@ -49,7 +49,7 @@ THE SOFTWARE.
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -59,7 +59,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
 {
     *ptr = tex1Dfetch<T>(textureObject, x);
@@ -67,7 +67,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -77,7 +77,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
 {
     *ptr = tex1D<T>(textureObject, x);
@@ -85,7 +85,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -95,7 +95,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
 {
     *ptr = tex2D<T>(textureObject, x, y);
@@ -103,7 +103,7 @@ static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -113,7 +113,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
 {
     *ptr = tex3D<T>(textureObject, x, y, z);
@@ -121,7 +121,7 @@ static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -131,7 +131,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
 {
     *ptr = tex1DLayered<T>(textureObject, x, layer);
@@ -139,7 +139,7 @@ static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t t
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -149,7 +149,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
 {
     *ptr = tex1DLayered<T>(textureObject, x, y, layer);
@@ -157,7 +157,7 @@ static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t t
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__  T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -167,7 +167,7 @@ static __device__ __hip_img_chk__  T texCubemap(hipTextureObject_t textureObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
 {
     *ptr = texCubemap<T>(textureObject, x, y, z);
@@ -175,7 +175,7 @@ static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t tex
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -185,7 +185,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
 {
     *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
@@ -193,7 +193,7 @@ static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObjec
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -224,7 +224,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
 {
     *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
@@ -232,7 +232,7 @@ static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t te
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -242,7 +242,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
 {
     *ptr = tex1DLod<T>(textureObject, x, level);
@@ -250,7 +250,7 @@ static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textu
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -260,7 +260,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
 {
     *ptr = tex2DLod<T>(textureObject, x, y, level);
@@ -268,7 +268,7 @@ static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textu
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -278,7 +278,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     *ptr = tex3DLod<T>(textureObject, x, y, z, level);
@@ -286,7 +286,7 @@ static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textu
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -296,7 +296,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
 {
     *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
@@ -304,7 +304,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__  T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -314,7 +314,7 @@ static __device__ __hip_img_chk__  T tex2DLayeredLod(hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
 {
     *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
@@ -322,7 +322,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -332,7 +332,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
 {
     *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
@@ -340,7 +340,7 @@ static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -352,7 +352,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
@@ -360,7 +360,7 @@ static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -370,7 +370,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
 {
     *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
@@ -378,7 +378,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureOb
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -388,7 +388,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
 {
     *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
@@ -396,7 +396,7 @@ static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t text
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -406,7 +406,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
 {
     *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
@@ -414,7 +414,7 @@ static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t text
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -424,7 +424,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
 {
     *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
@@ -432,7 +432,7 @@ static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t text
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -442,7 +442,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
 {
     *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
@@ -450,7 +450,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -460,7 +460,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
 {
     *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
@@ -468,7 +468,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__  T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
 {
     TEXTURE_OBJECT_PARAMETERS_INIT
@@ -480,7 +480,7 @@ static __device__ __hip_img_chk__  T texCubemapLayeredGrad(hipTextureObject_t te
 
 template <
     typename T,
-    typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
+    typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
 static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
 {
     *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
@@ -490,4 +490,4 @@ static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureO
 
 #if defined(__clang__)
 #pragma clang diagnostic pop
-#endif
+#endif

projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake

@@ -60,6 +60,13 @@ function(get_hiprtc_macros HIPRTC_DEFINES)
 #define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H\n\
 #define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H\n\
 #if !__HIP_NO_STD_DEFS__\n\
+#if defined(_WIN32)\n\
+typedef unsigned long long uint64_t;\n\
+typedef signed long long int64_t;\n\
+#else\n\
+typedef unsigned long uint64_t;\n\
+typedef signed long int64_t;\n\
+#endif\n\
 #if defined(__HIPRTC_PTRDIFF_T_IS_LONG_LONG__) && __HIPRTC_PTRDIFF_T_IS_LONG_LONG__==1\n\
 typedef long long ptrdiff_t;\n\
 #else\n\
@@ -67,6 +74,8 @@ typedef __PTRDIFF_TYPE__ ptrdiff_t;\n\
 #endif\n\
 typedef long clock_t;\n\
 namespace std {\n\
+using ::uint64_t;\n\
+using ::int64_t;\n\
 using ::ptrdiff_t;\n\
 using ::clock_t;\n\
 }\n\