AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
This reverts commit 353dbe6e3b.

Reason for revert: This is considered a breaking change and requires
multiple apps to change their behavior. This will be reintroduced in later releases.

Change-Id: I3481627115af1872785585a155cc6a0ecfbe1372
This commit is contained in:
Satyanvesh Dittakavi
2023-12-01 07:13:54 +00:00
committed by Maneesh Gupta
parent d347f1d67d
commit 7febad13be
9 changed files with 240 additions and 217 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_hip_fp16.h
+9 -9
View File
@@ -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<bool cond, typename T = void>
using Enable_if_t = typename __hip_internal::enable_if<cond, T>::type;
using Enable_if_t = typename std::enable_if<cond, T>::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<T>{}>* = nullptr>
Enable_if_t<std::is_floating_point<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<T>{}>* = nullptr>
typename T, Enable_if_t<std::is_integral<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<T>{}>* = nullptr>
Enable_if_t<std::is_floating_point<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<T>{}>* = nullptr>
typename T, Enable_if_t<std::is_integral<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<T>{}>* = nullptr>
Enable_if_t<std::is_floating_point<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<T>{}>* = nullptr>
typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
__HOST_DEVICE__
operator T() const { return data; }
#endif
hipamd/include/hip/amd_detail/amd_hip_runtime.h
+4
View File
@@ -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__)
hipamd/include/hip/amd_detail/amd_hip_vector_types.h
+135 -53
View File
@@ -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 <array>
#include <iosfwd>
#include <type_traits>
#else
namespace std {
using ::size_t;
template <class _Tp, _Tp __v> struct integral_constant {
static constexpr const _Tp value = __v;
typedef _Tp value_type;
typedef integral_constant type;
constexpr operator value_type() const { return value; }
constexpr value_type operator()() const { return value; }
};
template <class _Tp, _Tp __v> constexpr const _Tp integral_constant<_Tp, __v>::value;
typedef integral_constant<bool, true> true_type;
typedef integral_constant<bool, false> false_type;
template <bool B> using bool_constant = integral_constant<bool, B>;
typedef bool_constant<true> true_type;
typedef bool_constant<false> false_type;
template <bool __B, class __T = void> struct enable_if {};
template <class __T> struct enable_if<true, __T> { typedef __T type; };
template<bool _B> struct true_or_false_type : public false_type {};
template<> struct true_or_false_type<true> : public true_type {};
template <class _Tp> struct is_integral : public false_type {};
template <> struct is_integral<bool> : public true_type {};
template <> struct is_integral<char> : public true_type {};
template <> struct is_integral<signed char> : public true_type {};
template <> struct is_integral<unsigned char> : public true_type {};
template <> struct is_integral<wchar_t> : public true_type {};
template <> struct is_integral<short> : public true_type {};
template <> struct is_integral<unsigned short> : public true_type {};
template <> struct is_integral<int> : public true_type {};
template <> struct is_integral<unsigned int> : public true_type {};
template <> struct is_integral<long> : public true_type {};
template <> struct is_integral<unsigned long> : public true_type {};
template <> struct is_integral<long long> : public true_type {};
template <> struct is_integral<unsigned long long> : public true_type {};
template <class _Tp> struct is_arithmetic : public false_type {};
template <> struct is_arithmetic<bool> : public true_type {};
template <> struct is_arithmetic<char> : public true_type {};
template <> struct is_arithmetic<signed char> : public true_type {};
template <> struct is_arithmetic<unsigned char> : public true_type {};
template <> struct is_arithmetic<wchar_t> : public true_type {};
template <> struct is_arithmetic<short> : public true_type {};
template <> struct is_arithmetic<unsigned short> : public true_type {};
template <> struct is_arithmetic<int> : public true_type {};
template <> struct is_arithmetic<unsigned int> : public true_type {};
template <> struct is_arithmetic<long> : public true_type {};
template <> struct is_arithmetic<unsigned long> : public true_type {};
template <> struct is_arithmetic<long long> : public true_type {};
template <> struct is_arithmetic<unsigned long long> : public true_type {};
template <> struct is_arithmetic<float> : public true_type {};
template <> struct is_arithmetic<double> : public true_type {};
template<typename _Tp> struct is_floating_point : public false_type {};
template<> struct is_floating_point<float> : public true_type {};
template<> struct is_floating_point<double> : public true_type {};
template<> struct is_floating_point<long double> : public true_type {};
template <typename __T, typename __U> struct is_same : public false_type {};
template <typename __T> struct is_same<__T, __T> : public true_type {};
template<typename _Tp, bool = is_arithmetic<_Tp>::value>
struct is_signed : public false_type {};
template<typename _Tp>
struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {};
template <class _T1, class _T2> struct is_convertible
: public true_or_false_type<__is_convertible_to(_T1, _T2)> {};
template<typename _CharT> struct char_traits;
template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_istream;
template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_ostream;
typedef basic_istream<char> istream;
typedef basic_ostream<char> ostream;
template <typename __T> struct is_scalar : public integral_constant<bool, __is_scalar(__T)> {};
} // Namespace std.
#endif // defined(__HIPCC_RTC__)
namespace hip_impl {
@@ -194,7 +276,7 @@ THE SOFTWARE.
template<
typename U = T,
typename __hip_internal::enable_if<__hip_internal::is_signed<U>{}>::type* = nullptr>
typename std::enable_if<std::is_signed<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U, T>::value>::type* = nullptr>
typename std::enable_if<
std::is_convertible<U, T>::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<U, T>{}>::type* = nullptr>
typename std::enable_if<
std::is_convertible<U, T>{}>::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<U, T>{}>::type* = nullptr>
typename std::enable_if<
std::is_convertible<U, T>{}>::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<U, T>{}>::type* = nullptr>
typename std::enable_if<
std::is_convertible<U, T>{}>::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<U, T>{}>::type* = nullptr>
typename std::enable_if<
std::is_convertible<U, T>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_signed<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<U>{}>::type* = nullptr>
typename std::enable_if<std::is_integral<U>{}>::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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<T>{}>* = nullptr>
typename std::enable_if<std::is_integral<T>{}>* = 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<U>::value>::type,
typename __hip_internal::enable_if<__hip_internal::is_integral<T>{}>* = nullptr>
typename std::enable_if<std::is_arithmetic<U>::value>::type,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
__HOST_DEVICE__
inline
constexpr
@@ -1028,28 +1110,28 @@ THE SOFTWARE.
* Map HIP_vector_type<U, rankU> to HIP_vector_type<T, rankT>
*/
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 1 && rankU >= 1),
__forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 1 && rankU >= 1),
const HIP_vector_type<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT>(static_cast<T>(u.x));
};
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU == 1),
__forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU == 1),
const HIP_vector_type<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(0));
};
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 2 && rankU >= 2),
__forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 2 && rankU >= 2),
const HIP_vector_type<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(u.y));
};
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__forceinline__ __HOST_DEVICE__ typename __hip_internal::enable_if<(rankT == 4 && rankU == 1),
__forceinline__ __HOST_DEVICE__ typename std::enable_if<(rankT == 4 && rankU == 1),
const HIP_vector_type<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(0),
@@ -1057,7 +1139,7 @@ THE SOFTWARE.
};
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__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<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT>(static_cast<T>(u.x), static_cast<T>(u.y),
@@ -1065,7 +1147,7 @@ THE SOFTWARE.
};
template <typename T, unsigned int rankT, typename U, unsigned int rankU>
__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<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU>& u) {
return HIP_vector_type<T, rankT> (static_cast<T>(u.x), static_cast<T>(u.y),
hipamd/include/hip/amd_detail/amd_surface_functions.h
+14 -14
View File
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
int boundaryMode = hipBoundaryModeZero) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+4 -6
View File
@@ -62,20 +62,18 @@ namespace cooperative_groups {
/* Global scope */
template <unsigned int size>
using is_power_of_2 = __hip_internal::integral_constant<bool, (size & (size - 1)) == 0>;
using is_power_of_2 = std::integral_constant<bool, (size & (size - 1)) == 0>;
template <unsigned int size>
using is_valid_wavefront = __hip_internal::integral_constant<bool, (size <= __AMDGCN_WAVEFRONT_SIZE)>;
using is_valid_wavefront = std::integral_constant<bool, (size <= __AMDGCN_WAVEFRONT_SIZE)>;
template <unsigned int size>
using is_valid_tile_size =
__hip_internal::integral_constant<bool, is_power_of_2<size>::value &&
is_valid_wavefront<size>::value>;
std::integral_constant<bool, is_power_of_2<size>::value && is_valid_wavefront<size>::value>;
template <typename T>
using is_valid_type =
__hip_internal::integral_constant<bool, __hip_internal::is_integral<T>::value ||
__hip_internal::is_floating_point<T>::value>;
std::integral_constant<bool, std::is_integral<T>::value || std::is_floating_point<T>::value>;
namespace internal {
hipamd/include/hip/amd_detail/host_defines.h
+4 -56
View File
@@ -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<typename _Tp, bool = is_arithmetic<_Tp>::value>
template<typename _Tp>
struct is_signed<_Tp, true> : public true_or_false_type<_Tp(-1) < _Tp(0)> {};
template<class T>
auto test_returnable(int) -> decltype(
void(static_cast<T(*)()>(nullptr)), true_type{});
template<class>
auto test_returnable(...) -> false_type;
template<class T>
struct type_identity { using type = T; };
template<class T> // Note that `cv void&` is a substitution failure
auto try_add_lvalue_reference(int) -> type_identity<T&>;
template<class T> // Handle T = cv void case
auto try_add_lvalue_reference(...) -> type_identity<T>;
template<class T>
auto try_add_rvalue_reference(int) -> type_identity<T&&>;
template<class T>
auto try_add_rvalue_reference(...) -> type_identity<T>;
template<class T>
struct add_lvalue_reference
: decltype(try_add_lvalue_reference<T>(0)) {};
template<class T>
struct add_rvalue_reference
: decltype(try_add_rvalue_reference<T>(0)) {};
template<typename T>
typename add_rvalue_reference<T>::type declval() noexcept;
template<class From, class To>
auto test_implicitly_convertible(int) -> decltype(
void(declval<void(&)(To)>()(declval<From>())), true_type{});
template<class, class>
auto test_implicitly_convertible(...) -> false_type;
template<class T> struct remove_cv { typedef T type; };
template<class T> struct remove_cv<const T> { typedef T type; };
template<class T> struct remove_cv<volatile T> { typedef T type; };
template<class T> struct remove_cv<const volatile T> { typedef T type; };
template<class T>
struct is_void : public is_same<void, typename remove_cv<T>::type> {};
template<class From, class To>
struct is_convertible : public integral_constant<bool,
(decltype(test_returnable<To>(0))::value &&
decltype(test_implicitly_convertible<From, To>(0))::value) ||
(is_void<From>::value && is_void<To>::value)> {};
template<typename _CharT> struct char_traits;
template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_istream;
template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_ostream;
@@ -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
hipamd/include/hip/amd_detail/texture_fetch_functions.h
+24 -24
View File
@@ -39,13 +39,13 @@ template<typename T>
struct __hip_is_tex_surf_scalar_channel_type
{
static constexpr bool value =
__hip_internal::is_same<T, char>::value ||
__hip_internal::is_same<T, unsigned char>::value ||
__hip_internal::is_same<T, short>::value ||
__hip_internal::is_same<T, unsigned short>::value ||
__hip_internal::is_same<T, int>::value ||
__hip_internal::is_same<T, unsigned int>::value ||
__hip_internal::is_same<T, float>::value;
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<typename T>
@@ -71,10 +71,10 @@ template<typename T>
struct __hip_is_tex_normalized_channel_type
{
static constexpr bool value =
__hip_internal::is_same<T, char>::value ||
__hip_internal::is_same<T, unsigned char>::value ||
__hip_internal::is_same<T, short>::value ||
__hip_internal::is_same<T, unsigned short>::value;
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value;
};
template<
@@ -95,7 +95,7 @@ template <
typename Enable = void>
struct __hip_tex_ret
{
static_assert(__hip_internal::is_same<Enable, void>::value, "Invalid channel type!");
static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
};
/*
@@ -103,7 +103,7 @@ struct __hip_tex_ret
*/
template<typename T, typename U>
__forceinline__ __device__
typename __hip_internal::enable_if<
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<T>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(T) < sizeof(float)) {
@@ -126,7 +126,7 @@ __hipMapFrom(const U &u) {
*/
template<typename T, typename U>
__forceinline__ __device__
typename __hip_internal::enable_if<
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
@@ -149,7 +149,7 @@ __hipMapFrom(const U &u) {
*/
template<typename U, typename T>
__forceinline__ __device__
typename __hip_internal::enable_if<
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<T>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(T) < sizeof(float)) {
@@ -174,7 +174,7 @@ __hipMapTo(const T &t) {
*/
template<typename U, typename T>
__forceinline__ __device__
typename __hip_internal::enable_if<
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
@@ -203,7 +203,7 @@ template <typename T>
struct __hip_tex_ret<
T,
hipReadModeElementType,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
{
using type = T;
};
@@ -214,7 +214,7 @@ template<
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
};
@@ -223,7 +223,7 @@ template<typename T>
struct __hip_tex_ret<
T,
hipReadModeNormalizedFloat,
typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float;
};
@@ -234,7 +234,7 @@ template<
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeNormalizedFloat,
typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
};
@@ -426,7 +426,7 @@ template <
typename Enable = void>
struct __hip_tex2dgather_ret
{
static_assert(__hip_internal::is_same<Enable, void>::value, "Invalid channel type!");
static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
};
template <
@@ -438,7 +438,7 @@ template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeElementType,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
@@ -449,7 +449,7 @@ template<
struct __hip_tex2dgather_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
@@ -458,7 +458,7 @@ template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeNormalizedFloat,
typename __hip_internal::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float4;
};
hipamd/include/hip/amd_detail/texture_indirect_functions.h
+46 -46
View File
@@ -38,7 +38,7 @@ THE SOFTWARE.
template <
typename T,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -48,7 +48,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
template <
typename T,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
{
*ptr = tex1Dfetch<T>(textureObject, x);
@@ -56,7 +56,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex
template <
typename T,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -66,7 +66,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa
template <
typename T,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
{
*ptr = tex1D<T>(textureObject, x);
@@ -74,7 +74,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO
template <
typename T,
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
{
*ptr = tex2D<T>(textureObject, x, y);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = tex3D<T>(textureObject, x, y, z);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, layer);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, y, layer);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = texCubemap<T>(textureObject, x, y, z);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, comp);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
{
*ptr = tex1DLod<T>(textureObject, x, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
{
*ptr = tex2DLod<T>(textureObject, x, y, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = tex3DLod<T>(textureObject, x, y, z, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
{
*ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
*ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = texCubemapLod<T>(textureObject, x, y, z, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
*ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
*ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
*ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
*ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
*ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -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<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
hipamd/src/hiprtc/cmake/HIPRTC.cmake
-9
View File
@@ -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\