Move type traits to CK namespace

include/ck/tensor_operation/gpu/device/tensor_layout.hpp

@@ -406,7 +406,7 @@ struct G_NDHW : public BaseTensorLayout
 
 template <
     typename Layout,
-    typename std::enable_if<std::is_base_of<BaseTensorLayout, Layout>::value, bool>::type = false>
+    typename ck::enable_if<ck::is_base_of<BaseTensorLayout, Layout>::value, bool>::type = false>
 std::ostream& operator<<(std::ostream& os, const Layout&)
 {
     os << Layout::name;

include/ck/utility/amd_wave_read_first_lane.hpp

@@ -100,8 +100,8 @@ __device__ inline int32_t amd_wave_read_first_lane(int32_t value)
 }
 
 template <typename Object,
-          typename = ck::enable_if_t<std::is_class<Object>::value &&
-                                     std::is_trivially_copyable<Object>::value>>
+          typename = ck::enable_if_t<ck::is_class<Object>::value &&
+                                     ck::is_trivially_copyable<Object>::value>>
 __device__ auto amd_wave_read_first_lane(const Object& obj)
 {
     using Size                = unsigned;

include/ck/utility/array.hpp

@@ -52,7 +52,7 @@ template <typename X, typename... Xs>
 __host__ __device__ constexpr auto make_array(X&& x, Xs&&... xs)
 {
     using data_type = remove_cvref_t<X>;
-    return Array<data_type, sizeof...(Xs) + 1>{std::forward<X>(x), std::forward<Xs>(xs)...};
+    return Array<data_type, sizeof...(Xs) + 1>{ck::forward<X>(x), ck::forward<Xs>(xs)...};
 }
 
 // make empty array

include/ck/utility/container_helper.hpp

@@ -326,14 +326,14 @@ template <typename T, index_t NX, index_t NY>
 __host__ __device__ constexpr auto container_concat(const Array<T, NX>& ax, const Array<T, NY>& ay)
 {
     return unpack2(
-        [&](auto&&... zs) { return make_array(std::forward<decltype(zs)>(zs)...); }, ax, ay);
+        [&](auto&&... zs) { return make_array(ck::forward<decltype(zs)>(zs)...); }, ax, ay);
 }
 
 template <typename... X, typename... Y>
 __host__ __device__ constexpr auto container_concat(const Tuple<X...>& tx, const Tuple<Y...>& ty)
 {
     return unpack2(
-        [&](auto&&... zs) { return make_tuple(std::forward<decltype(zs)>(zs)...); }, tx, ty);
+        [&](auto&&... zs) { return make_tuple(ck::forward<decltype(zs)>(zs)...); }, tx, ty);
 }
 
 template <typename Container>

include/ck/utility/functional.hpp

@@ -120,11 +120,11 @@ constexpr auto conditional_expr(X&& x, Y&& y)
 {
     if constexpr(predicate)
     {
-        return std::forward<X>(x);
+        return ck::forward<X>(x);
     }
     else
     {
-        return std::forward<Y>(y);
+        return ck::forward<Y>(y);
     }
 }
 

include/ck/utility/functional4.hpp

@@ -21,7 +21,7 @@ struct unpack_impl<Sequence<Is...>>
     template <typename F, typename X>
     __host__ __device__ constexpr auto operator()(F&& f, X&& x) const
     {
-        return std::forward<F>(f)(std::forward<X>(x).At(Number<Is>{})...);
+        return ck::forward<F>(f)(ck::forward<X>(x).At(Number<Is>{})...);
     }
 };
 
@@ -35,8 +35,8 @@ struct unpack2_impl<Sequence<Is...>, Sequence<Js...>>
     template <typename F, typename X, typename Y>
     __host__ __device__ constexpr auto operator()(F&& f, X&& x, Y&& y) const
     {
-        return std::forward<F>(f)(std::forward<X>(x).At(Number<Is>{})...,
-                                  std::forward<Y>(y).At(Number<Js>{})...);
+        return ck::forward<F>(f)(ck::forward<X>(x).At(Number<Is>{})...,
+                                 ck::forward<Y>(y).At(Number<Js>{})...);
     }
 };
 
@@ -47,7 +47,7 @@ __host__ __device__ constexpr auto unpack(F&& f, X&& x)
 {
     using X_ = remove_reference_t<X>;
     return detail::unpack_impl<typename arithmetic_sequence_gen<0, X_::Size(), 1>::type>{}(
-        std::forward<F>(f), std::forward<X>(x));
+        ck::forward<F>(f), ck::forward<X>(x));
 }
 
 // TODO: properly implement unpack that takes any number of containers
@@ -58,7 +58,7 @@ __host__ __device__ constexpr auto unpack2(F&& f, X&& x, Y&& y)
     using Y_ = remove_reference_t<Y>;
     return detail::unpack2_impl<typename arithmetic_sequence_gen<0, X_::Size(), 1>::type,
                                 typename arithmetic_sequence_gen<0, Y_::Size(), 1>::type>{}(
-        std::forward<F>(f), std::forward<X>(x), std::forward<Y>(y));
+        ck::forward<F>(f), ck::forward<X>(x), ck::forward<Y>(y));
 }
 
 } // namespace ck

include/ck/utility/integral_constant.hpp

@@ -48,4 +48,9 @@ __host__ __device__ constexpr auto operator%(integral_constant<TX, X>, integral_
     return integral_constant<decltype(X % Y), X % Y>{};
 }
 
+template <bool B>
+using bool_constant = integral_constant<bool, B>;
+
+using true_type  = bool_constant<true>;
+using false_type = bool_constant<false>;
 } // namespace ck

include/ck/utility/tuple.hpp

@@ -32,7 +32,7 @@ struct TupleElementKeyData
     template <typename T,
               typename enable_if<!is_same<remove_cvref_t<T>, TupleElementKeyData>::value,
                                  bool>::type = false>
-    __host__ __device__ constexpr TupleElementKeyData(T&& v) : mData(std::forward<T>(v))
+    __host__ __device__ constexpr TupleElementKeyData(T&& v) : mData(ck::forward<T>(v))
     {
     }
 
@@ -67,7 +67,7 @@ get_tuple_element_data_reference(TupleElementKeyData<Key, Data>&& x)
 template <typename Key, typename Data>
 __host__ __device__ constexpr Data get_tuple_element_data(const TupleElementKeyData<Key, Data>& x)
 {
-    return std::forward(x.mData);
+    return ck::forward(x.mData);
 }
 
 template <typename Indices, typename... Xs>
@@ -83,13 +83,13 @@ struct TupleImpl<Sequence<Is...>, Xs...> : TupleElementKeyData<TupleElementKey<I
                                      !is_same<remove_cvref_t<Y>, TupleImpl>::value,
                                  bool>::type = false>
     __host__ __device__ constexpr TupleImpl(Y&& y)
-        : TupleElementKeyData<TupleElementKey<Is>, Xs>(std::forward<Y>(y))...
+        : TupleElementKeyData<TupleElementKey<Is>, Xs>(ck::forward<Y>(y))...
     {
     }
 
     template <typename... Ys, typename enable_if<sizeof...(Ys) >= 2, bool>::type = false>
     __host__ __device__ constexpr TupleImpl(Ys&&... ys)
-        : TupleElementKeyData<TupleElementKey<Is>, Xs>(std::forward<Ys>(ys))...
+        : TupleElementKeyData<TupleElementKey<Is>, Xs>(ck::forward<Ys>(ys))...
     {
         static_assert(sizeof...(Is) == sizeof...(Xs) && sizeof...(Is) == sizeof...(Ys),
                       "wrong! inconsistent size");
@@ -123,14 +123,14 @@ struct Tuple : detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(X
     template <typename Y,
               typename enable_if<sizeof...(Xs) == 1 && !is_same<remove_cvref_t<Y>, Tuple>::value,
                                  bool>::type = false>
-    __host__ __device__ constexpr Tuple(Y&& y) : base(std::forward<Y>(y))
+    __host__ __device__ constexpr Tuple(Y&& y) : base(ck::forward<Y>(y))
     {
     }
 
     template <typename... Ys,
               typename enable_if<sizeof...(Ys) == sizeof...(Xs) && sizeof...(Ys) >= 2, bool>::type =
                   false>
-    __host__ __device__ constexpr Tuple(Ys&&... ys) : base(std::forward<Ys>(ys)...)
+    __host__ __device__ constexpr Tuple(Ys&&... ys) : base(ck::forward<Ys>(ys)...)
     {
     }
 
@@ -208,7 +208,7 @@ using tuple_element_t = typename tuple_element<I, TTuple>::type;
 template <typename... Xs>
 __host__ __device__ constexpr auto make_tuple(Xs&&... xs)
 {
-    return Tuple<remove_cvref_t<Xs>...>(std::forward<Xs>(xs)...);
+    return Tuple<remove_cvref_t<Xs>...>(ck::forward<Xs>(xs)...);
 }
 
 // https://en.cppreference.com/w/cpp/utility/tuple/tie

include/ck/utility/tuple_helper.hpp

@@ -28,7 +28,7 @@ __host__ __device__ constexpr auto concat_tuple_of_reference(const Tuple<X&...>&
                                                              const Tuple<Y&...>& ty)
 {
     return unpack2(
-        [&](auto&&... zs) { return Tuple<decltype(zs)...>{std::forward<decltype(zs)>(zs)...}; },
+        [&](auto&&... zs) { return Tuple<decltype(zs)...>{ck::forward<decltype(zs)>(zs)...}; },
         tx,
         ty);
 }

include/ck/utility/type.hpp

@@ -7,8 +7,8 @@
 #include "ck/utility/enable_if.hpp"
 #include "ck/utility/integral_constant.hpp"
 
+namespace ck {
 #ifdef __HIPCC_RTC__
-namespace std {
 // NOLINTNEXTLINE
 #define CK_BUILTIN_TYPE_TRAIT1(name)         \
     template <class T>                       \
@@ -31,55 +31,42 @@ namespace std {
     }
 
 CK_BUILTIN_TYPE_TRAIT1(is_class);
+CK_BUILTIN_TYPE_TRAIT1(is_pointer);
 CK_BUILTIN_TYPE_TRAIT1(is_reference);
 CK_BUILTIN_TYPE_TRAIT1(is_trivially_copyable);
 CK_BUILTIN_TYPE_TRAIT1(is_unsigned);
 CK_BUILTIN_TYPE_TRAIT2(is_base_of);
-CK_BUILTIN_TYPE_TRAITN(is_constructible);
-CK_BUILTIN_TYPE_TRAITN(is_nothrow_constructible);
 
 template <class T>
-struct remove_reference
-{
-    typedef T type;
-};
-template <class T>
-struct remove_reference<T&>
-{
-    typedef T type;
-};
-template <class T>
-struct remove_reference<T&&>
+struct remove_cv
 {
-    typedef T type;
+    using type = T;
 };
 
 template <class T>
-struct remove_const
+struct remove_cv<const T> : remove_cv<T>
 {
-    typedef T type;
 };
+
 template <class T>
-struct remove_const<const T>
+struct remove_cv<volatile T> : remove_cv<T>
 {
-    typedef T type;
 };
 
 template <class T>
-struct remove_volatile
+struct remove_reference
 {
     typedef T type;
 };
 template <class T>
-struct remove_volatile<volatile T>
+struct remove_reference<T&>
 {
     typedef T type;
 };
-
 template <class T>
-struct remove_cv
+struct remove_reference<T&&>
 {
-    typedef typename remove_volatile<typename remove_const<T>::type>::type type;
+    typedef T type;
 };
 
 template <class T>
@@ -108,21 +95,6 @@ struct remove_pointer<T* const volatile>
     typedef T type;
 };
 
-template <class T>
-struct is_pointer_helper : std::false_type
-{
-};
-
-template <class T>
-struct is_pointer_helper<T*> : std::true_type
-{
-};
-
-template <class T>
-struct is_pointer : is_pointer_helper<typename std::remove_cv<T>::type>
-{
-};
-
 template <typename T>
 constexpr T&& forward(typename remove_reference<T>::type& t_) noexcept
 {
@@ -134,12 +106,20 @@ constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
 {
     return static_cast<T&&>(t_);
 }
-
-template <typename T>
-inline constexpr bool is_reference_v = is_reference<T>::value;
-} // namespace std
+#else
+#include <utility>
+#include <type_traits>
+using std::forward;
+using std::is_base_of;
+using std::is_class;
+using std::is_pointer;
+using std::is_reference;
+using std::is_trivially_copyable;
+using std::is_unsigned;
+using std::remove_cv;
+using std::remove_pointer;
+using std::remove_reference;
 #endif
-namespace ck {
 
 template <typename X, typename Y>
 struct is_same : public integral_constant<bool, false>
@@ -151,28 +131,39 @@ struct is_same<X, X> : public integral_constant<bool, true>
 {
 };
 
+template <typename T>
+inline constexpr bool is_reference_v = is_reference<T>::value;
+
 template <typename X, typename Y>
 inline constexpr bool is_same_v = is_same<X, Y>::value;
 
+template <typename X, typename Y>
+inline constexpr bool is_base_of_v = is_base_of<X, Y>::value;
+
+template <typename T>
+inline constexpr bool is_unsigned_v = is_unsigned<T>::value;
+
 template <typename T>
-using remove_reference_t = typename std::remove_reference<T>::type;
+using remove_reference_t = typename remove_reference<T>::type;
 
 template <typename T>
-using remove_reference_t = typename std::remove_reference<T>::type;
+using remove_reference_t = typename remove_reference<T>::type;
 
 template <typename T>
-using remove_cv_t = typename std::remove_cv<T>::type;
+using remove_cv_t = typename remove_cv<T>::type;
 
 template <typename T>
 using remove_cvref_t = remove_cv_t<remove_reference_t<T>>;
 
 template <typename T>
-using remove_pointer_t = typename std::remove_pointer<T>::type;
+using remove_pointer_t = typename remove_pointer<T>::type;
 
 template <typename T>
-inline constexpr bool is_pointer_v = std::is_pointer<T>::value;
+inline constexpr bool is_pointer_v = is_pointer<T>::value;
 
-template <typename Y, typename X, typename enable_if<sizeof(X) == sizeof(Y), bool>::type = false>
+template <typename Y,
+          typename X,
+          typename ck::enable_if<sizeof(X) == sizeof(Y), bool>::type = false>
 __host__ __device__ constexpr Y bit_cast(const X& x)
 {
 #if CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST

include/ck/utility/type_convert.hpp

@@ -13,7 +13,7 @@ namespace ck {
 template <typename Y, typename X>
 __host__ __device__ constexpr Y type_convert(X x)
 {
-    static_assert(!std::is_reference_v<Y> && !std::is_reference_v<X>);
+    static_assert(!ck::is_reference_v<Y> && !ck::is_reference_v<X>);
 
     return static_cast<Y>(x);
 }