Skip to content

[ESIMD] Fix specialization/instantiation order in traits infra. #6677

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Aug 31, 2022
Merged

[ESIMD] Fix specialization/instantiation order in traits infra. #6677

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 3 additions & 21 deletions sycl/include/sycl/ext/intel/esimd/detail/elem_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -153,14 +153,12 @@ struct element_type_traits<T, std::enable_if_t<is_vectorizable_v<T>>> {

// ------------------- Useful meta-functions and declarations

template <class T> using __raw_t = typename element_type_traits<T>::RawT;
template <class T>
using __raw_t = typename __ESIMD_DNS::element_type_traits<T>::RawT;
template <class T>
using __cpp_t = typename __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
using __cpp_t = typename element_type_traits<T>::EnclosingCppT;

template <class T, int N>
using __raw_vec_t =
vector_type_t<typename __ESIMD_DNS::element_type_traits<T>::RawT, N>;
using __raw_vec_t = vector_type_t<typename element_type_traits<T>::RawT, N>;

// Note: using RawVecT in comparison result type calculation does *not* mean
// the comparison is actually performed on the raw types.
Expand Down Expand Up @@ -602,22 +600,6 @@ vector_comparison_op_traits<Op, WrapperT, N>::impl(__raw_vec_t<WrapperT, N> X,
vector_comparison_op_default<Op, T1, N>(X1, Y1));
}

// Proxy class to access bit representation of a wrapper type both on host and
// device. Declared as friend to the wrapper types (e.g. sycl::half).
// Specific type traits implementations (scalar_conversion_traits) can use
// concrete wrapper type specializations of the static functions in this class
// to access private fields in the wrapper type (e.g. sycl::half).
// TODO add this functionality to sycl type implementation? With C++20,
// std::bit_cast should be a good replacement.
class WrapperElementTypeProxy {
public:
template <class WrapperT>
static inline __raw_t<WrapperT> bitcast_to_raw_scalar(WrapperT Val);

template <class WrapperT>
static inline WrapperT bitcast_to_wrapper_scalar(__raw_t<WrapperT> Val);
};

// "Generic" version of std::is_floating_point_v which returns "true" also for
// the wrapper floating-point types such as sycl::half.
template <typename T>
Expand Down
79 changes: 44 additions & 35 deletions sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,22 +20,28 @@ namespace sycl {
__SYCL_INLINE_VER_NAMESPACE(_V1) {
namespace ext::intel::esimd::detail {

template <class T>
struct element_type_traits<T, std::enable_if_t<std::is_same_v<T, sycl::half>>> {
// Can't use sycl::detail::half_impl::StorageT as RawT for both host and
// device as it still maps to struct on/ host (even though the struct is a
// trivial wrapper around uint16_t), and for ESIMD we need a type which can be
// an element of clang vector.
// Standalone definitions to use w/o instantiating element_type_traits.
#ifdef __SYCL_DEVICE_ONLY__
// Can't use sycl::detail::half_impl::StorageT as RawT for both host and
// device as it still maps to struct on/ host (even though the struct is a
// trivial wrapper around uint16_t), and for ESIMD we need a type which can be
// an element of clang vector.
using half_raw_type = sycl::detail::half_impl::StorageT;
// On device, _Float16 is native Cpp type, so it is the enclosing C++ type
using half_enclosing_cpp_type = half_raw_type;
#else
using half_raw_type = uint16_t;
using half_enclosing_cpp_type = float;
#endif // __SYCL_DEVICE_ONLY__

template <> struct element_type_traits<sycl::half> {
using RawT = half_raw_type;
using EnclosingCppT = half_enclosing_cpp_type;
#ifdef __SYCL_DEVICE_ONLY__
using RawT = sycl::detail::half_impl::StorageT;
// On device, _Float16 is native Cpp type, so it is the enclosing C++ type
using EnclosingCppT = RawT;
// On device, operations on half are translated to operations on _Float16,
// which is natively supported by the device compiler
static inline constexpr bool use_native_cpp_ops = true;
#else
using RawT = uint16_t;
using EnclosingCppT = float;
// On host, we can't use native Cpp '+', '-' etc. over uint16_t to emulate the
// operations on half type.
static inline constexpr bool use_native_cpp_ops = false;
Expand All @@ -47,8 +53,8 @@ struct element_type_traits<T, std::enable_if_t<std::is_same_v<T, sycl::half>>> {
// ------------------- Type conversion traits

template <int N> struct vector_conversion_traits<sycl::half, N> {
using StdT = __cpp_t<sycl::half>;
using RawT = __raw_t<sycl::half>;
using StdT = half_enclosing_cpp_type;
using RawT = half_raw_type;

static ESIMD_INLINE vector_type_t<RawT, N>
convert_to_raw(vector_type_t<StdT, N> Val)
Expand All @@ -57,7 +63,7 @@ template <int N> struct vector_conversion_traits<sycl::half, N> {
;
#else
{
vector_type_t<__raw_t<sycl::half>, N> Output = 0;
vector_type_t<half_raw_type, N> Output = 0;

for (int i = 0; i < N; i += 1) {
// 1. Convert Val[i] to float (x) using c++ static_cast
Expand Down Expand Up @@ -89,46 +95,49 @@ template <int N> struct vector_conversion_traits<sycl::half, N> {
#endif // __SYCL_DEVICE_ONLY__
};

// WrapperElementTypeProxy (a friend of sycl::half) must be used to access
// private fields of the sycl::half.
template <>
ESIMD_INLINE __raw_t<sycl::half>
WrapperElementTypeProxy::bitcast_to_raw_scalar<sycl::half>(sycl::half Val) {
// Proxy class to access bit representation of a wrapper type both on host and
// device. Declared as friend to the wrapper types (e.g. sycl::half).
// Specific type traits implementations (scalar_conversion_traits) can use
// concrete wrapper type specializations of the static functions in this class
// to access private fields in the wrapper type (e.g. sycl::half).
// TODO add this functionality to sycl type implementation? With C++20,
// std::bit_cast should be a good replacement.
class WrapperElementTypeProxy {
public:
static ESIMD_INLINE half_raw_type bitcast_to_raw_scalar(sycl::half Val) {
#ifdef __SYCL_DEVICE_ONLY__
return Val.Data;
return Val.Data;
#else
return Val.Data.Buf;
return Val.Data.Buf;
#endif // __SYCL_DEVICE_ONLY__
}
}

template <>
ESIMD_INLINE sycl::half
WrapperElementTypeProxy::bitcast_to_wrapper_scalar<sycl::half>(
__raw_t<sycl::half> Val) {
static ESIMD_INLINE sycl::half bitcast_to_wrapper_scalar(half_raw_type Val) {
#ifndef __SYCL_DEVICE_ONLY__
return sycl::half(::sycl::detail::host_half_impl::half(Val));
return sycl::half(::sycl::detail::host_half_impl::half(Val));
#else
sycl::half Res;
Res.Data = Val;
return Res;
sycl::half Res;
Res.Data = Val;
return Res;
#endif // __SYCL_DEVICE_ONLY__
}
}
};

template <> struct scalar_conversion_traits<sycl::half> {
using RawT = __raw_t<sycl::half>;
using RawT = half_raw_type;

static ESIMD_INLINE RawT bitcast_to_raw(sycl::half Val) {
return WrapperElementTypeProxy::bitcast_to_raw_scalar<sycl::half>(Val);
return WrapperElementTypeProxy::bitcast_to_raw_scalar(Val);
}

static ESIMD_INLINE sycl::half bitcast_to_wrapper(RawT Val) {
return WrapperElementTypeProxy::bitcast_to_wrapper_scalar<sycl::half>(Val);
return WrapperElementTypeProxy::bitcast_to_wrapper_scalar(Val);
}
};

#ifdef __SYCL_DEVICE_ONLY__
template <>
struct is_esimd_arithmetic_type<__raw_t<sycl::half>, void> : std::true_type {};
struct is_esimd_arithmetic_type<half_raw_type, void> : std::true_type {};
#endif // __SYCL_DEVICE_ONLY__

// Misc
Expand Down