diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 5ff2662570..62d88a3756 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -70,13 +70,13 @@ THE SOFTWARE. #include "amd_device_functions.h" #include "amd_warp_functions.h" #endif - namespace __hip_internal + namespace std { - template<> struct is_floating_point<_Float16> : __hip_internal::true_type {}; + template<> struct is_floating_point<_Float16> : std::true_type {}; } template - using Enable_if_t = typename __hip_internal::enable_if::type; + using Enable_if_t = typename std::enable_if::type; // BEGIN STRUCT __HALF struct __half { @@ -98,7 +98,7 @@ THE SOFTWARE. __half(decltype(data) x) : data{x} {} template< typename T, - Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> + Enable_if_t{}>* = nullptr> __HOST_DEVICE__ __half(T x) : data{static_cast<_Float16>(x)} {} #endif @@ -112,7 +112,7 @@ THE SOFTWARE. // CREATORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> + typename T, Enable_if_t{}>* = nullptr> __HOST_DEVICE__ __half(T x) : data{static_cast<_Float16>(x)} {} #endif @@ -157,7 +157,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< typename T, - Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> + Enable_if_t{}>* = nullptr> __HOST_DEVICE__ __half& operator=(T x) { @@ -169,7 +169,7 @@ THE SOFTWARE. // MANIPULATORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> + typename T, Enable_if_t{}>* = nullptr> __device__ __half& operator=(T x) { @@ -227,7 +227,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< typename T, - Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> + Enable_if_t{}>* = nullptr> __HOST_DEVICE__ operator T() const { return data; } #endif @@ -241,7 +241,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> + typename T, Enable_if_t{}>* = nullptr> __HOST_DEVICE__ operator T() const { return data; } #endif diff --git a/hipamd/include/hip/amd_detail/amd_hip_runtime.h b/hipamd/include/hip/amd_detail/amd_hip_runtime.h index 903de73fb1..f4a143798f 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_runtime.h +++ b/hipamd/include/hip/amd_detail/amd_hip_runtime.h @@ -90,10 +90,14 @@ 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__) 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 22be76316a..c4736e4743 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 - 2023 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2022 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,6 +49,88 @@ THE SOFTWARE. #include #include #include +#else +namespace std { +using ::size_t; + +template 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 constexpr const _Tp integral_constant<_Tp, __v>::value; + +typedef integral_constant true_type; +typedef integral_constant false_type; + +template using bool_constant = integral_constant; +typedef bool_constant true_type; +typedef bool_constant false_type; + +template struct enable_if {}; +template struct enable_if { typedef __T type; }; + +template struct true_or_false_type : public false_type {}; +template<> struct true_or_false_type : public true_type {}; + +template struct is_integral : public false_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; +template <> struct is_integral : public true_type {}; + +template struct is_arithmetic : public false_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; +template <> struct is_arithmetic : public true_type {}; + +template struct is_floating_point : public false_type {}; +template<> struct is_floating_point : public true_type {}; +template<> struct is_floating_point : public true_type {}; +template<> struct is_floating_point : public true_type {}; + +template struct is_same : public false_type {}; +template struct is_same<__T, __T> : public true_type {}; + +template::value> + struct is_signed : public false_type {}; +template + struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {}; + +template struct is_convertible + : public true_or_false_type<__is_convertible_to(_T1, _T2)> {}; + +template struct char_traits; +template> class basic_istream; +template> class basic_ostream; +typedef basic_istream istream; +typedef basic_ostream ostream; + +template struct is_scalar : public integral_constant {}; +} // Namespace std. #endif // defined(__HIPCC_RTC__) namespace hip_impl { @@ -194,7 +276,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_signed{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_ operator-() const noexcept { @@ -205,7 +287,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_ operator~() const noexcept { @@ -215,7 +297,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator%=(const Native_vec_& x_) noexcept { @@ -224,7 +306,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator^=(const Native_vec_& x_) noexcept { @@ -233,7 +315,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator|=(const Native_vec_& x_) noexcept { @@ -242,7 +324,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator&=(const Native_vec_& x_) noexcept { @@ -251,7 +333,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator>>=(const Native_vec_& x_) noexcept { @@ -260,7 +342,7 @@ THE SOFTWARE. } template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator<<=(const Native_vec_& x_) noexcept { @@ -367,8 +449,8 @@ THE SOFTWARE. HIP_vector_type() = default; template< typename U, - typename __hip_internal::enable_if< - __hip_internal::is_convertible::value>::type* = nullptr> + typename std::enable_if< + std::is_convertible::value>::type* = nullptr> __HOST_DEVICE__ explicit constexpr @@ -377,7 +459,7 @@ THE SOFTWARE. {} template< // TODO: constrain based on type as well. typename... Us, - typename __hip_internal::enable_if< + typename std::enable_if< (rank > 1) && sizeof...(Us) == rank>::type* = nullptr> __HOST_DEVICE__ constexpr @@ -437,8 +519,8 @@ THE SOFTWARE. } template< typename U, - typename __hip_internal::enable_if< - __hip_internal::is_convertible{}>::type* = nullptr> + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator+=(U x) noexcept { @@ -457,8 +539,8 @@ THE SOFTWARE. } template< typename U, - typename __hip_internal::enable_if< - __hip_internal::is_convertible{}>::type* = nullptr> + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator-=(U x) noexcept { @@ -484,8 +566,8 @@ THE SOFTWARE. template< typename U, - typename __hip_internal::enable_if< - __hip_internal::is_convertible{}>::type* = nullptr> + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator*=(U x) noexcept { @@ -510,8 +592,8 @@ THE SOFTWARE. } template< typename U, - typename __hip_internal::enable_if< - __hip_internal::is_convertible{}>::type* = nullptr> + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator/=(U x) noexcept { @@ -520,7 +602,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_signed{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type operator-() const noexcept { @@ -535,7 +617,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type operator~() const noexcept { @@ -550,7 +632,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept { @@ -564,7 +646,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept { @@ -578,7 +660,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept { @@ -592,7 +674,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept { @@ -606,7 +688,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept { @@ -620,7 +702,7 @@ THE SOFTWARE. template< typename U = T, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> + typename std::enable_if{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept { @@ -792,7 +874,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -805,7 +887,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -818,7 +900,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -831,7 +913,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -844,7 +926,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -857,7 +939,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -870,7 +952,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -883,7 +965,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -896,7 +978,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -909,7 +991,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -922,7 +1004,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -935,7 +1017,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -948,7 +1030,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -961,7 +1043,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -974,7 +1056,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -987,7 +1069,7 @@ THE SOFTWARE. template< typename T, unsigned int n, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1000,7 +1082,7 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1013,8 +1095,8 @@ THE SOFTWARE. typename T, unsigned int n, typename U, - typename __hip_internal::enable_if<__hip_internal::is_arithmetic::value>::type, - typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> + typename std::enable_if::value>::type, + typename std::enable_if{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1028,28 +1110,28 @@ THE SOFTWARE. * Map HIP_vector_type to HIP_vector_type */ template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 1 && rankU >= 1), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 1 && rankU >= 1), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type(static_cast(u.x)); }; template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU == 1), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU == 1), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(0)); }; template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU >= 2), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU >= 2), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(u.y)); }; template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 1), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 1), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(0), @@ -1057,7 +1139,7 @@ THE SOFTWARE. }; template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 2), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 2), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type(static_cast(u.x), static_cast(u.y), @@ -1065,7 +1147,7 @@ THE SOFTWARE. }; template - __forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 4), + __forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 4), const HIP_vector_type>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(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 108d67c4e5..556988a8e7 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -98,7 +98,7 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -109,7 +109,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -119,7 +119,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -129,7 +129,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -139,7 +139,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -149,7 +149,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -159,7 +159,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -169,7 +169,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -179,7 +179,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -189,7 +189,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -199,7 +199,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -209,7 +209,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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)); @@ -219,7 +219,7 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -230,7 +230,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 15607b9683..f62246a031 100644 --- a/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -62,20 +62,18 @@ namespace cooperative_groups { /* Global scope */ template -using is_power_of_2 = __hip_internal::integral_constant; +using is_power_of_2 = std::integral_constant; template -using is_valid_wavefront = __hip_internal::integral_constant; +using is_valid_wavefront = std::integral_constant; template using is_valid_tile_size = - __hip_internal::integral_constant::value && - is_valid_wavefront::value>; + std::integral_constant::value && is_valid_wavefront::value>; template using is_valid_type = - __hip_internal::integral_constant::value || - __hip_internal::is_floating_point::value>; + std::integral_constant::value || std::is_floating_point::value>; namespace internal { diff --git a/hipamd/include/hip/amd_detail/host_defines.h b/hipamd/include/hip/amd_detail/host_defines.h index c37f5d273f..0fad2b4704 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 - 2023 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2022 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,7 +33,8 @@ THE SOFTWARE. #define GENERIC_GRID_LAUNCH 1 #endif -#if defined(__cplusplus) +#if defined(__clang__) && defined(__HIP__) + namespace __hip_internal { typedef unsigned char uint8_t; typedef unsigned short uint16_t; @@ -111,57 +112,6 @@ template::value> template struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {}; -template - auto test_returnable(int) -> decltype( - void(static_cast(nullptr)), true_type{}); -template - auto test_returnable(...) -> false_type; - -template - struct type_identity { using type = T; }; - -template // Note that `cv void&` is a substitution failure - auto try_add_lvalue_reference(int) -> type_identity; -template // Handle T = cv void case - auto try_add_lvalue_reference(...) -> type_identity; - -template - auto try_add_rvalue_reference(int) -> type_identity; -template - auto try_add_rvalue_reference(...) -> type_identity; - -template -struct add_lvalue_reference - : decltype(try_add_lvalue_reference(0)) {}; - -template -struct add_rvalue_reference - : decltype(try_add_rvalue_reference(0)) {}; - -template -typename add_rvalue_reference::type declval() noexcept; - -template - auto test_implicitly_convertible(int) -> decltype( - void(declval()(declval())), true_type{}); - -template - auto test_implicitly_convertible(...) -> false_type; - -template struct remove_cv { typedef T type; }; -template struct remove_cv { typedef T type; }; -template struct remove_cv { typedef T type; }; -template struct remove_cv { typedef T type; }; - -template -struct is_void : public is_same::type> {}; - -template -struct is_convertible : public integral_constant(0))::value && - decltype(test_implicitly_convertible(0))::value) || - (is_void::value && is_void::value)> {}; - template struct char_traits; template> class basic_istream; template> class basic_ostream; @@ -186,9 +136,7 @@ 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)) @@ -227,6 +175,6 @@ typedef __hip_internal::int64_t __hip_int64_t; #define __constant__ #define __hip_img_chk__ -#endif // defined(__clang__) && defined(__HIP__) +#endif #endif diff --git a/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/hipamd/include/hip/amd_detail/texture_fetch_functions.h index 0e95b7a7c5..c4dcbe78d8 100644 --- a/hipamd/include/hip/amd_detail/texture_fetch_functions.h +++ b/hipamd/include/hip/amd_detail/texture_fetch_functions.h @@ -39,13 +39,13 @@ template struct __hip_is_tex_surf_scalar_channel_type { static constexpr bool value = - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value; + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value; }; template @@ -71,10 +71,10 @@ template struct __hip_is_tex_normalized_channel_type { static constexpr bool value = - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value || - __hip_internal::is_same::value; + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value; }; template< @@ -95,7 +95,7 @@ template < typename Enable = void> struct __hip_tex_ret { - static_assert(__hip_internal::is_same::value, "Invalid channel type!"); + static_assert(std::is_same::value, "Invalid channel type!"); }; /* @@ -103,7 +103,7 @@ struct __hip_tex_ret */ template __forceinline__ __device__ -typename __hip_internal::enable_if< +typename std::enable_if< __hip_is_tex_surf_scalar_channel_type::value, const T>::type __hipMapFrom(const U &u) { if constexpr (sizeof(T) < sizeof(float)) { @@ -126,7 +126,7 @@ __hipMapFrom(const U &u) { */ template __forceinline__ __device__ -typename __hip_internal::enable_if< +typename std::enable_if< __hip_is_tex_surf_scalar_channel_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 __forceinline__ __device__ -typename __hip_internal::enable_if< +typename std::enable_if< __hip_is_tex_surf_scalar_channel_type::value, const U>::type __hipMapTo(const T &t) { if constexpr (sizeof(T) < sizeof(float)) { @@ -174,7 +174,7 @@ __hipMapTo(const T &t) { */ template __forceinline__ __device__ -typename __hip_internal::enable_if< +typename std::enable_if< __hip_is_tex_surf_scalar_channel_type::value, const U>::type __hipMapTo(const T &t) { if constexpr (sizeof(typename T::value_type) < sizeof(float)) { @@ -203,7 +203,7 @@ template struct __hip_tex_ret< T, hipReadModeElementType, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = T; }; @@ -214,7 +214,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeElementType, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -223,7 +223,7 @@ template struct __hip_tex_ret< T, hipReadModeNormalizedFloat, - typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> { using type = float; }; @@ -234,7 +234,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeNormalizedFloat, - typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -426,7 +426,7 @@ template < typename Enable = void> struct __hip_tex2dgather_ret { - static_assert(__hip_internal::is_same::value, "Invalid channel type!"); + static_assert(std::is_same::value, "Invalid channel type!"); }; template < @@ -438,7 +438,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeElementType, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = HIP_vector_type; }; @@ -449,7 +449,7 @@ template< struct __hip_tex2dgather_ret< HIP_vector_type, hipReadModeElementType, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type; }; @@ -458,7 +458,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeNormalizedFloat, - typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type::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 24e1e6bb98..1e435018df 100644 --- a/hipamd/include/hip/amd_detail/texture_indirect_functions.h +++ b/hipamd/include/hip/amd_detail/texture_indirect_functions.h @@ -38,7 +38,7 @@ THE SOFTWARE. template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x) { *ptr = tex1Dfetch(textureObject, x); @@ -56,7 +56,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x) { *ptr = tex1D(textureObject, x); @@ -74,7 +74,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -84,7 +84,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y) { *ptr = tex2D(textureObject, x, y); @@ -92,7 +92,7 @@ static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -102,7 +102,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = tex3D(textureObject, x, y, z); @@ -110,7 +110,7 @@ static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -120,7 +120,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer) { *ptr = tex1DLayered(textureObject, x, layer); @@ -128,7 +128,7 @@ static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -138,7 +138,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer) { *ptr = tex1DLayered(textureObject, x, y, layer); @@ -146,7 +146,7 @@ static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -156,7 +156,7 @@ static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = texCubemap(textureObject, x, y, z); @@ -164,7 +164,7 @@ static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t tex template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -174,7 +174,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, layer); @@ -182,7 +182,7 @@ static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObjec template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -213,7 +213,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0) { *ptr = texCubemapLayered(textureObject, x, y, comp); @@ -221,7 +221,7 @@ static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t te template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -231,7 +231,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level) { *ptr = tex1DLod(textureObject, x, level); @@ -239,7 +239,7 @@ static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -249,7 +249,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level) { *ptr = tex2DLod(textureObject, x, y, level); @@ -257,7 +257,7 @@ static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -267,7 +267,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, level); @@ -275,7 +275,7 @@ static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -285,7 +285,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level) { *ptr = tex1DLayeredLod(textureObject, x, layer, level); @@ -293,7 +293,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_ template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -303,7 +303,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, layer, level); @@ -311,7 +311,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_ template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -321,7 +321,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, level); @@ -329,7 +329,7 @@ static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -341,7 +341,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, dPdx, dPdy); @@ -349,7 +349,7 @@ static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -359,7 +359,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, layer, level); @@ -367,7 +367,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureOb template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -377,7 +377,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { *ptr = tex1DGrad(textureObject, x, dPdx, dPdy); @@ -385,7 +385,7 @@ static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -395,7 +395,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, dPdx, dPdy); @@ -403,7 +403,7 @@ static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -413,7 +413,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, dPdx, dPdy); @@ -421,7 +421,7 @@ static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -431,7 +431,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, layer, dPdx, dPdy); @@ -439,7 +439,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -449,7 +449,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, layer, dPdx, dPdy); @@ -457,7 +457,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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 @@ -469,7 +469,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t te template < typename T, - typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::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(textureObject, x, y, z, layer, dPdx, dPdy); diff --git a/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/hipamd/src/hiprtc/cmake/HIPRTC.cmake index 680e4da429..ef9e559d1f 100644 --- a/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -60,13 +60,6 @@ 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\ @@ -74,8 +67,6 @@ 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\