simd.hpp

Go to the documentation of this file.

//==------------ - simd.hpp - DPC++ Explicit SIMD API   --------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// Implement Explicit SIMD vector APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/ext/intel/esimd/detail/simd_mask_impl.hpp>
#include <sycl/ext/intel/esimd/detail/simd_obj_impl.hpp>
 
#include <sycl/ext/intel/esimd/detail/intrin.hpp>
#include <sycl/ext/intel/esimd/detail/memory_intrin.hpp>
#include <sycl/ext/intel/esimd/detail/sycl_util.hpp>
#include <sycl/ext/intel/esimd/detail/types.hpp>
#include <sycl/ext/intel/esimd/simd_view.hpp>
 
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
 
#ifndef __SYCL_DEVICE_ONLY__
#include <sycl/detail/iostream_proxy.hpp>
#endif // __SYCL_DEVICE_ONLY__
 
namespace sycl {
__SYCL_INLINE_VER_NAMESPACE(_V1) {
namespace ext::intel::esimd {
 
 
 
 
 
template <typename Ty, int N>
class simd : public detail::simd_obj_impl<
                 detail::__raw_t<Ty>, N, simd<Ty, N>,
                 std::enable_if_t<detail::is_valid_simd_elem_type_v<Ty>>> {
  using base_type = detail::simd_obj_impl<detail::__raw_t<Ty>, N, simd<Ty, N>>;
 
public:
  using base_type::base_type;
  using element_type = Ty;
  using raw_element_type = typename base_type::raw_element_type;
  using raw_vector_type = typename base_type::raw_vector_type;
  static constexpr int length = N;
 
  template <typename SimdT,
            class = std::enable_if_t<__ESIMD_DNS::is_simd_type_v<SimdT> &&
                                     (length == SimdT::length)>>
  simd(const SimdT &RHS)
      : base_type(detail::convert_vector<Ty, detail::element_type_t<SimdT>, N>(
            RHS.data())) {
    __esimd_dbg_print(simd(const SimdT &RHS));
  }
 
  // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
  template <
      int N1 = N, class Ty1 = Ty,
      class SFINAE = std::enable_if_t<
          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
  simd(const sycl::ext::oneapi::experimental::simd<Ty, N1> &v)
      : simd(static_cast<raw_vector_type>(v)) {}
 
  template <typename T1,
            class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
  simd(T1 Val) : base_type(Val) {
    __esimd_dbg_print(simd(T1 Val));
  }
 
  template <class To, class T = simd,
            class = std::enable_if_t<(T::length == 1) &&
                                     detail::is_valid_simd_elem_type_v<To>>>
  operator To() const {
    __esimd_dbg_print(operator To());
    return detail::convert_scalar<To, element_type>(base_type::data()[0]);
  }
 
  template <
      int N1, class Ty1 = Ty,
      class SFINAE = std::enable_if_t<
          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
  operator sycl::ext::oneapi::experimental::simd<Ty, N1>() {
    return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
  }
 
  simd &operator++() {
    *this += 1;
    return *this;
  }
 
  simd operator++(int) {
    simd Ret(base_type::data());
    operator++();
    return Ret;
  }
 
  simd &operator--() {
    *this -= 1;
    return *this;
  }
 
  simd operator--(int) {
    simd Ret(base_type::data());
    operator--();
    return Ret;
  }
 
#define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID)                    \
  template <class T1 = Ty> simd operator ARITH_UNARY_OP() const {              \
    static_assert(!std::is_unsigned_v<T1>,                                     \
                  #ARITH_UNARY_OP "doesn't apply to unsigned types");          \
    return simd{detail::vector_unary_op<detail::UnaryOp::ID, T1, N>(           \
        base_type::data())};                                                   \
  }
 
  __ESIMD_DEF_SIMD_ARITH_UNARY_OP(-, minus)
  __ESIMD_DEF_SIMD_ARITH_UNARY_OP(+, plus)
#undef __ESIMD_DEF_SIMD_ARITH_UNARY_OP
};
 
 
 
template <typename To, typename From, int N>
ESIMD_INLINE simd<To, N> convert(const simd<From, N> &val) {
  if constexpr (std::is_same_v<std::remove_const_t<To>,
                               std::remove_const_t<From>>)
    return val;
  else
    return detail::convert_vector<To, From, N>(val.data());
}
 
 
template <int N> using simd_mask = detail::simd_mask_type<N>;
 
 
} // namespace ext::intel::esimd
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
} // namespace sycl
 
template <typename Ty, int N>
std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
#ifdef __SYCL_DEVICE_ONLY__
    {}
#else
{
  OS << "{";
  for (int I = 0; I < N; I++) {
    OS << V[I];
    if (I < N - 1)
      OS << ",";
  }
  OS << "}";
  return OS;
}
#endif // __SYCL_DEVICE_ONLY__