From b27e2abc41f63b948b70f6cf3ce227b874344744 Mon Sep 17 00:00:00 2001
From: "Dittakavi, Satyanvesh" <Satyanvesh.Dittakavi@amd.com>
Date: Thu, 22 May 2025 13:27:19 +0530
Subject: [PATCH] SWDEV-482125 - HIPRTC implementation of standard headers
 moved from std to __hip_internal namespace (#376)

Change-Id: Ia359ed45e0916ad397401113a69a6b26a590d807
---
 hipamd/include/hip/amd_detail/amd_hip_fp16.h  |  20 +-
 .../hip/amd_detail/amd_hip_vector_types.h     | 206 ++++++------------
 .../hip/amd_detail/amd_surface_functions.h    |  30 +--
 .../hip_cooperative_groups_helper.h           |  10 +-
 hipamd/include/hip/amd_detail/host_defines.h  |  60 ++++-
 .../hip/amd_detail/texture_fetch_functions.h  |  50 ++---
 .../amd_detail/texture_indirect_functions.h   |  94 ++++----
 7 files changed, 224 insertions(+), 246 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/hipamd/include/hip/amd_detail/amd_hip_fp16.h
index b39a4f3755..d8ba155081 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_fp16.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_fp16.h
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2025 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
@@ -78,13 +78,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 {
@@ -106,7 +106,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
@@ -120,7 +120,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
@@ -165,7 +165,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)
                 {
@@ -177,7 +177,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)
                 {
@@ -235,7 +235,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
@@ -249,7 +249,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
diff --git a/hipamd/include/hip/amd_detail/amd_hip_vector_types.h b/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
index 9b472c534b..27d33561b9 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
+++ b/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 - 2025 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
@@ -49,88 +49,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 {
@@ -263,7 +181,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
             {
@@ -274,7 +192,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
             {
@@ -284,7 +202,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
             {
@@ -293,7 +211,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
             {
@@ -302,7 +220,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
             {
@@ -311,7 +229,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
             {
@@ -320,7 +238,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
             {
@@ -329,7 +247,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
             {
@@ -410,12 +328,12 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
       __HOST_DEVICE__
       HIP_vector_type() = default;
       template <typename U,
-                typename std::enable_if<std::is_convertible<U, T>::value>::type* = nullptr>
+                typename __hip_internal::enable_if<std::is_convertible<U, T>::value>::type* = nullptr>
       __HOST_DEVICE__ explicit constexpr HIP_vector_type(U x_) noexcept
           : HIP_vector_base<T, rank>{static_cast<T>(x_)} {}
       template <  // TODO: constrain based on type as well.
           typename... Us,
-          typename std::enable_if<(rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
+          typename __hip_internal::enable_if<(rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
       __HOST_DEVICE__ constexpr HIP_vector_type(Us... xs) noexcept
           : HIP_vector_base<T, rank>{static_cast<T>(xs)...} {}
       __HOST_DEVICE__
@@ -455,8 +373,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return tmp;
       }
 
-      __HOST_DEVICE__
-      HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept {
+      __HOST_DEVICE__ HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) += get_native_vector(x);
 #else
@@ -464,13 +381,13 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 #endif
         return *this;
       }
-      template <typename U, typename std::enable_if<std::is_convertible<U, T>{}>::type* = nullptr>
+      template <typename U, typename __hip_internal::enable_if<
+                    __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator+=(U x) noexcept {
         return *this += make_vector_type<T, rank>(x);
       }
 
-      __HOST_DEVICE__
-      HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept {
+      __HOST_DEVICE__ HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) -= get_native_vector(x);
 #else
@@ -478,13 +395,13 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 #endif
         return *this;
       }
-      template <typename U, typename std::enable_if<std::is_convertible<U, T>{}>::type* = nullptr>
+      template <typename U, typename __hip_internal::enable_if<
+                    __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator-=(U x) noexcept {
         return *this -= make_vector_type<T, rank>(x);
       }
 
-      __HOST_DEVICE__
-      HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept {
+      __HOST_DEVICE__ HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) *= get_native_vector(x);
 #else
@@ -498,7 +415,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return HIP_vector_type{x} *= y;
       }
 
-      template <typename U, typename std::enable_if<std::is_convertible<U, T>{}>::type* = nullptr>
+      template <typename U, typename __hip_internal::enable_if<
+                    __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator*=(U x) noexcept {
         return *this *= make_vector_type<T, rank>(x);
       }
@@ -508,8 +426,7 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return HIP_vector_type{x} /= y;
       }
 
-      __HOST_DEVICE__
-      HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept {
+      __HOST_DEVICE__ HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) /= get_native_vector(x);
 #else
@@ -517,12 +434,14 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
 #endif
         return *this;
       }
-      template <typename U, typename std::enable_if<std::is_convertible<U, T>{}>::type* = nullptr>
+      template <typename U, typename __hip_internal::enable_if<
+                    __hip_internal::is_convertible<U, T>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator/=(U x) noexcept {
         return *this /= make_vector_type<T, rank>(x);
       }
 
-      template <typename U = T, typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
+      template <typename U = T, typename __hip_internal::enable_if<
+                    __hip_internal::is_signed<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type operator-() const noexcept {
         auto tmp(*this);
 #if __HIP_USE_NATIVE_VECTOR__
@@ -533,7 +452,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return tmp;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type operator~() const noexcept {
         HIP_vector_type r{*this};
 #if __HIP_USE_NATIVE_VECTOR__
@@ -544,8 +464,9 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return r;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-      __HOST_DEVICE__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept {
+      template<typename U = T, typename __hip_internal::enable_if<
+                    __hip_internal::is_integral<U>{}>::type* = nullptr>
+       __HOST_DEVICE__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) %= get_native_vector(x);
 #else
@@ -554,7 +475,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return *this;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) ^= get_native_vector(x);
@@ -564,7 +486,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return *this;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) |= get_native_vector(x);
@@ -574,8 +497,9 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return *this;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-      __HOST_DEVICE__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept {
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
+       __HOST_DEVICE__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) &= get_native_vector(x);
 #else
@@ -584,7 +508,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return *this;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) >>= get_native_vector(x);
@@ -594,7 +519,8 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
         return *this;
       }
 
-      template <typename U = T, typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+      template<typename U = T, typename __hip_internal::enable_if<
+                   __hip_internal::is_integral<U>{}>::type* = nullptr>
       __HOST_DEVICE__ HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept {
 #if __HIP_USE_NATIVE_VECTOR__
         get_native_vector(*this) <<= get_native_vector(x);
@@ -761,7 +687,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
@@ -774,7 +700,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
@@ -787,7 +713,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
@@ -800,7 +726,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
@@ -813,7 +739,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
@@ -826,7 +752,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
@@ -839,7 +765,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
@@ -852,7 +778,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
@@ -865,7 +791,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
@@ -878,7 +804,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
@@ -891,7 +817,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
@@ -904,7 +830,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
@@ -917,7 +843,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
@@ -930,7 +856,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
@@ -943,7 +869,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
@@ -956,7 +882,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
@@ -969,7 +895,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
@@ -982,8 +908,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
@@ -997,28 +923,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),
@@ -1026,7 +952,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),
@@ -1034,7 +960,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),
diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h
index c931cf62d1..3f62011945 100644
--- a/hipamd/include/hip/amd_detail/amd_surface_functions.h
+++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2018 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2018 - 2025 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
@@ -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
@@ -129,7 +129,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));
@@ -148,7 +148,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
     int2 coords{x, y};
@@ -168,7 +168,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
     int2 coords{x, y};
@@ -189,7 +189,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
     int4 coords{x, y, z, 0};
@@ -210,7 +210,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
     int4 coords{x, y, z, 0};
@@ -230,7 +230,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));
@@ -249,7 +249,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));
@@ -269,7 +269,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
     int2 coords{x, y};
@@ -290,7 +290,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
     int2 coords{x, y};
@@ -311,7 +311,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
     int2 coords{x, y};
@@ -332,7 +332,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
     int2 coords{x, y};
@@ -354,7 +354,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
@@ -377,7 +377,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
diff --git a/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
index 1c0e6227ad..0afae29da9 100644
--- a/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+++ b/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2025 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
@@ -66,18 +66,18 @@ 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 <= 64>;
+using is_valid_wavefront = __hip_internal::integral_constant<bool, size <= 64>;
 
 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 {
 
diff --git a/hipamd/include/hip/amd_detail/host_defines.h b/hipamd/include/hip/amd_detail/host_defines.h
index e7e8364969..e0a3c1f23d 100644
--- a/hipamd/include/hip/amd_detail/host_defines.h
+++ b/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 - 2025 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;
@@ -140,7 +190,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))
@@ -179,6 +231,6 @@ typedef __hip_internal::int64_t __hip_int64_t;
 #define __constant__
 
 #define __hip_img_chk__
-#endif
+#endif  // defined(__clang__) && defined(__HIP__)
 
 #endif
diff --git a/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/hipamd/include/hip/amd_detail/texture_fetch_functions.h
index 4de7f5ebd9..4123ea130b 100644
--- a/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+++ b/hipamd/include/hip/amd_detail/texture_fetch_functions.h
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2025 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
@@ -39,13 +39,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>
@@ -71,10 +71,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<
@@ -95,7 +95,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!");
 };
 
 /*
@@ -103,7 +103,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)) {
@@ -126,7 +126,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)) {
@@ -149,7 +149,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)) {
@@ -174,7 +174,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)) {
@@ -203,7 +203,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;
 };
@@ -214,7 +214,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>;
 };
@@ -223,7 +223,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;
 };
@@ -234,7 +234,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>;
 };
@@ -449,7 +449,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 <
@@ -461,7 +461,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>;
 };
@@ -472,7 +472,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>;
 };
@@ -481,7 +481,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;
 };
diff --git a/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/hipamd/include/hip/amd_detail/texture_indirect_functions.h
index bbc46ccdff..6d7a328b6b 100644
--- a/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+++ b/hipamd/include/hip/amd_detail/texture_indirect_functions.h
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - 2025 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
@@ -38,7 +38,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
@@ -48,7 +48,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);
@@ -56,7 +56,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
@@ -66,7 +66,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);
@@ -74,7 +74,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
@@ -85,7 +85,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);
@@ -93,7 +93,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
@@ -104,7 +104,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);
@@ -112,7 +112,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
@@ -123,7 +123,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);
@@ -131,7 +131,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
@@ -142,7 +142,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);
@@ -150,7 +150,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
@@ -161,7 +161,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);
@@ -169,7 +169,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
@@ -180,7 +180,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);
@@ -188,7 +188,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
@@ -220,7 +220,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);
@@ -228,7 +228,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
@@ -238,7 +238,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);
@@ -246,7 +246,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
@@ -257,7 +257,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);
@@ -265,7 +265,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
@@ -276,7 +276,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);
@@ -284,7 +284,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
@@ -295,7 +295,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);
@@ -303,7 +303,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
@@ -333,7 +333,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);
@@ -341,7 +341,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
@@ -354,7 +354,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);
@@ -362,7 +362,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
@@ -373,7 +373,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);
@@ -381,7 +381,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
@@ -391,7 +391,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);
@@ -399,7 +399,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
@@ -411,7 +411,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);
@@ -419,7 +419,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
@@ -433,7 +433,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);
@@ -441,7 +441,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
@@ -452,7 +452,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);
@@ -460,7 +460,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
@@ -472,7 +472,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);
@@ -480,7 +480,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
@@ -493,7 +493,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);