From 239fdc3bfd8dca7c09d465995755a7e564c231fc Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Mon, 9 Oct 2023 17:44:37 +0000 Subject: [PATCH] SWDEV-419480 - HIPRTC implementation of standard headers shouldn't be in std namespace HIPRTC implementation of few type traits are defined in std namespace which causes conflicts with the std C++ headers when apps try to include those std headers. Instead, use the specific implementation of these headers from __hip_internal namespace to avoid this. Change-Id: Ia358df159af8783d1eba9a19cd458e192bf7e9e6 [ROCm/clr commit: 353dbe6e3b592555d6d0128a28b500c39d74746d] --- .../include/hip/amd_detail/amd_hip_fp16.h | 18 +- .../include/hip/amd_detail/amd_hip_runtime.h | 4 - .../hip/amd_detail/amd_hip_vector_types.h | 188 +++++------------- .../hip/amd_detail/amd_surface_functions.h | 28 +-- .../hip_cooperative_groups_helper.h | 10 +- .../include/hip/amd_detail/host_defines.h | 60 +++++- .../hip/amd_detail/texture_fetch_functions.h | 48 ++--- .../amd_detail/texture_indirect_functions.h | 94 ++++----- .../clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake | 9 + 9 files changed, 218 insertions(+), 241 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 1fa2ba0d4c..5124afbe39 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -84,13 +84,13 @@ THE SOFTWARE. #include "amd_device_functions.h" #include "amd_warp_functions.h" #endif - namespace std + namespace __hip_internal { - template<> struct is_floating_point<_Float16> : std::true_type {}; + template<> struct is_floating_point<_Float16> : __hip_internal::true_type {}; } template - using Enable_if_t = typename std::enable_if::type; + using Enable_if_t = typename __hip_internal::enable_if::type; // BEGIN STRUCT __HALF struct __half { @@ -112,7 +112,7 @@ THE SOFTWARE. __half(decltype(data) x) : data{x} {} template< typename T, - Enable_if_t{}>* = nullptr> + Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> __HOST_DEVICE__ __half(T x) : data{static_cast<_Float16>(x)} {} #endif @@ -126,7 +126,7 @@ THE SOFTWARE. // CREATORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t{}>* = nullptr> + typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ __half(T x) : data{static_cast<_Float16>(x)} {} #endif @@ -171,7 +171,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< typename T, - Enable_if_t{}>* = nullptr> + Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> __HOST_DEVICE__ __half& operator=(T x) { @@ -183,7 +183,7 @@ THE SOFTWARE. // MANIPULATORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t{}>* = nullptr> + typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> __device__ __half& operator=(T x) { @@ -241,7 +241,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< typename T, - Enable_if_t{}>* = nullptr> + Enable_if_t<__hip_internal::is_floating_point{}>* = nullptr> __HOST_DEVICE__ operator T() const { return data; } #endif @@ -255,7 +255,7 @@ THE SOFTWARE. #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< - typename T, Enable_if_t{}>* = nullptr> + typename T, Enable_if_t<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ operator T() const { return data; } #endif diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h index f4a143798f..903de73fb1 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h @@ -90,14 +90,10 @@ size_t amd_dbgapi_get_build_id(); #else #if !__HIP_NO_STD_DEFS__ typedef unsigned int uint32_t; -typedef unsigned long long uint64_t; typedef signed int int32_t; -typedef signed long long int64_t; namespace std { using ::uint32_t; -using ::uint64_t; using ::int32_t; -using ::int64_t; } #endif // __HIP_NO_STD_DEFS__ #endif // !defined(__HIPCC_RTC__) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h index f6ffc1d5ed..b643d247f1 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -48,88 +48,6 @@ THE SOFTWARE. #include #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 { @@ -275,7 +193,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_signed{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_ operator-() const noexcept { @@ -286,7 +204,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_ operator~() const noexcept { @@ -296,7 +214,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator%=(const Native_vec_& x_) noexcept { @@ -305,7 +223,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator^=(const Native_vec_& x_) noexcept { @@ -314,7 +232,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator|=(const Native_vec_& x_) noexcept { @@ -323,7 +241,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator&=(const Native_vec_& x_) noexcept { @@ -332,7 +250,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator>>=(const Native_vec_& x_) noexcept { @@ -341,7 +259,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ Native_vec_& operator<<=(const Native_vec_& x_) noexcept { @@ -448,8 +366,8 @@ template struct is_scalar : public integral_constant::value>::type* = nullptr> + typename __hip_internal::enable_if< + __hip_internal::is_convertible::value>::type* = nullptr> __HOST_DEVICE__ explicit constexpr @@ -458,7 +376,7 @@ template struct is_scalar : public integral_constant 1) && sizeof...(Us) == rank>::type* = nullptr> __HOST_DEVICE__ constexpr @@ -514,8 +432,8 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if< + __hip_internal::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator+=(U x) noexcept { @@ -530,8 +448,8 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if< + __hip_internal::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator-=(U x) noexcept { @@ -553,8 +471,8 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if< + __hip_internal::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator*=(U x) noexcept { @@ -575,8 +493,8 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if< + __hip_internal::is_convertible{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator/=(U x) noexcept { @@ -585,7 +503,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_signed{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type operator-() const noexcept { @@ -596,7 +514,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type operator~() const noexcept { @@ -607,7 +525,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept { @@ -617,7 +535,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept { @@ -627,7 +545,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept { @@ -637,7 +555,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept { @@ -647,7 +565,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept { @@ -657,7 +575,7 @@ template struct is_scalar : public integral_constant{}>::type* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>::type* = nullptr> __HOST_DEVICE__ HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept { @@ -825,7 +743,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -838,7 +756,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -851,7 +769,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -864,7 +782,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -877,7 +795,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -890,7 +808,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -903,7 +821,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -916,7 +834,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -929,7 +847,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -942,7 +860,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -955,7 +873,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -968,7 +886,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -981,7 +899,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -994,7 +912,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1007,7 +925,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1020,7 +938,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1033,7 +951,7 @@ template struct is_scalar : public integral_constant{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1046,8 +964,8 @@ template struct is_scalar : public integral_constant::value>::type, - typename std::enable_if{}>* = nullptr> + typename __hip_internal::enable_if<__hip_internal::is_arithmetic::value>::type, + typename __hip_internal::enable_if<__hip_internal::is_integral{}>* = nullptr> __HOST_DEVICE__ inline constexpr @@ -1061,28 +979,28 @@ template struct is_scalar : public integral_constant to HIP_vector_type */ template - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type(static_cast(u.x)); }; template - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(0)); }; template - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(u.y)); }; template - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(0), @@ -1090,7 +1008,7 @@ template struct is_scalar : public integral_constant - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type(static_cast(u.x), static_cast(u.y), @@ -1098,7 +1016,7 @@ template struct is_scalar : public integral_constant - __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>::type __hipMapVector(const HIP_vector_type& u) { return HIP_vector_type (static_cast(u.x), static_cast(u.y), diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h index 5974fb6374..b14e3d6520 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -111,7 +111,7 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -122,7 +122,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -132,7 +132,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -142,7 +142,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -152,7 +152,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -162,7 +162,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -172,7 +172,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -182,7 +182,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -192,7 +192,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -202,7 +202,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -212,7 +212,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -222,7 +222,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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)); @@ -232,7 +232,7 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -243,7 +243,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index a451b144ea..358cbe14ad 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -70,18 +70,20 @@ namespace cooperative_groups { /* Global scope */ template -using is_power_of_2 = std::integral_constant; +using is_power_of_2 = __hip_internal::integral_constant; template -using is_valid_wavefront = std::integral_constant; +using is_valid_wavefront = __hip_internal::integral_constant; template using is_valid_tile_size = - std::integral_constant::value && is_valid_wavefront::value>; + __hip_internal::integral_constant::value && + is_valid_wavefront::value>; template using is_valid_type = - std::integral_constant::value || std::is_floating_point::value>; + __hip_internal::integral_constant::value || + __hip_internal::is_floating_point::value>; namespace internal { diff --git a/projects/clr/hipamd/include/hip/amd_detail/host_defines.h b/projects/clr/hipamd/include/hip/amd_detail/host_defines.h index 0fad2b4704..c37f5d273f 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/host_defines.h +++ b/projects/clr/hipamd/include/hip/amd_detail/host_defines.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -33,8 +33,7 @@ THE SOFTWARE. #define GENERIC_GRID_LAUNCH 1 #endif -#if defined(__clang__) && defined(__HIP__) - +#if defined(__cplusplus) namespace __hip_internal { typedef unsigned char uint8_t; typedef unsigned short uint16_t; @@ -112,6 +111,57 @@ template::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; @@ -136,7 +186,9 @@ typedef __hip_internal::int8_t __hip_int8_t; typedef __hip_internal::int16_t __hip_int16_t; typedef __hip_internal::int32_t __hip_int32_t; typedef __hip_internal::int64_t __hip_int64_t; +#endif // defined(__cplusplus) +#if defined(__clang__) && defined(__HIP__) #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __host__ __attribute__((host)) #define __device__ __attribute__((device)) @@ -175,6 +227,6 @@ typedef __hip_internal::int64_t __hip_int64_t; #define __constant__ #define __hip_img_chk__ -#endif +#endif // defined(__clang__) && defined(__HIP__) #endif diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h index baf2b28039..dc45561ff6 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h @@ -51,13 +51,13 @@ template struct __hip_is_tex_surf_scalar_channel_type { static constexpr bool 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; + __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; }; template @@ -83,10 +83,10 @@ template struct __hip_is_tex_normalized_channel_type { static constexpr bool value = - std::is_same::value || - std::is_same::value || - std::is_same::value || - std::is_same::value; + __hip_internal::is_same::value || + __hip_internal::is_same::value || + __hip_internal::is_same::value || + __hip_internal::is_same::value; }; template< @@ -107,7 +107,7 @@ template < typename Enable = void> struct __hip_tex_ret { - static_assert(std::is_same::value, "Invalid channel type!"); + static_assert(__hip_internal::is_same::value, "Invalid channel type!"); }; /* @@ -115,7 +115,7 @@ struct __hip_tex_ret */ template __forceinline__ __device__ -typename std::enable_if< +typename __hip_internal::enable_if< __hip_is_tex_surf_scalar_channel_type::value, const T>::type __hipMapFrom(const U &u) { if constexpr (sizeof(T) < sizeof(float)) { @@ -138,7 +138,7 @@ __hipMapFrom(const U &u) { */ template __forceinline__ __device__ -typename std::enable_if< +typename __hip_internal::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)) { @@ -161,7 +161,7 @@ __hipMapFrom(const U &u) { */ template __forceinline__ __device__ -typename std::enable_if< +typename __hip_internal::enable_if< __hip_is_tex_surf_scalar_channel_type::value, const U>::type __hipMapTo(const T &t) { if constexpr (sizeof(T) < sizeof(float)) { @@ -186,7 +186,7 @@ __hipMapTo(const T &t) { */ template __forceinline__ __device__ -typename std::enable_if< +typename __hip_internal::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)) { @@ -215,7 +215,7 @@ template struct __hip_tex_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = T; }; @@ -226,7 +226,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -235,7 +235,7 @@ template struct __hip_tex_ret< T, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> { using type = float; }; @@ -246,7 +246,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_tex_normalized_channel_type>::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -438,7 +438,7 @@ template < typename Enable = void> struct __hip_tex2dgather_ret { - static_assert(std::is_same::value, "Invalid channel type!"); + static_assert(__hip_internal::is_same::value, "Invalid channel type!"); }; template < @@ -450,7 +450,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = HIP_vector_type; }; @@ -461,7 +461,7 @@ template< struct __hip_tex2dgather_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type; }; @@ -470,7 +470,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> + typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> { using type = float4; }; diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h index 54e77723e8..fb8d2a5347 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h @@ -49,7 +49,7 @@ THE SOFTWARE. template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -59,7 +59,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -67,7 +67,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -77,7 +77,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -85,7 +85,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -95,7 +95,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -103,7 +103,7 @@ static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -113,7 +113,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -121,7 +121,7 @@ static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -131,7 +131,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -139,7 +139,7 @@ static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -149,7 +149,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -157,7 +157,7 @@ static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -167,7 +167,7 @@ static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -175,7 +175,7 @@ static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t tex template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -185,7 +185,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -193,7 +193,7 @@ static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObjec template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -224,7 +224,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -232,7 +232,7 @@ static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t te template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -242,7 +242,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -250,7 +250,7 @@ static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -260,7 +260,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -268,7 +268,7 @@ static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -278,7 +278,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -286,7 +286,7 @@ static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -296,7 +296,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -304,7 +304,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_ template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -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::value>::type* = nullptr> + typename __hip_internal::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); @@ -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::value>::type* = nullptr> + typename __hip_internal::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 @@ -332,7 +332,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -340,7 +340,7 @@ static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -352,7 +352,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -360,7 +360,7 @@ static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -370,7 +370,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -378,7 +378,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureOb template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -388,7 +388,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -396,7 +396,7 @@ static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -406,7 +406,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -414,7 +414,7 @@ static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -424,7 +424,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -432,7 +432,7 @@ static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -442,7 +442,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -450,7 +450,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -460,7 +460,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -468,7 +468,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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 @@ -480,7 +480,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t te template < typename T, - typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> + typename __hip_internal::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); @@ -490,4 +490,4 @@ static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureO #if defined(__clang__) #pragma clang diagnostic pop -#endif \ No newline at end of file +#endif diff --git a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index bfd09e8fe7..6bb7c83487 100644 --- a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -60,6 +60,13 @@ function(get_hiprtc_macros HIPRTC_DEFINES) #define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H\n\ #define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H\n\ #if !__HIP_NO_STD_DEFS__\n\ +#if defined(_WIN32)\n\ +typedef unsigned long long uint64_t;\n\ +typedef signed long long int64_t;\n\ +#else\n\ +typedef unsigned long uint64_t;\n\ +typedef signed long int64_t;\n\ +#endif\n\ #if defined(__HIPRTC_PTRDIFF_T_IS_LONG_LONG__) && __HIPRTC_PTRDIFF_T_IS_LONG_LONG__==1\n\ typedef long long ptrdiff_t;\n\ #else\n\ @@ -67,6 +74,8 @@ typedef __PTRDIFF_TYPE__ ptrdiff_t;\n\ #endif\n\ typedef long clock_t;\n\ namespace std {\n\ +using ::uint64_t;\n\ +using ::int64_t;\n\ using ::ptrdiff_t;\n\ using ::clock_t;\n\ }\n\