DPC++ Runtime
Runtime libraries for oneAPI DPC++
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE > Class Template Reference

This is a base class for all ESIMD simd classes with real storage (simd, simd_mask_impl). More...

#include <sycl/ext/intel/esimd/detail/simd_obj_impl.hpp>

Collaboration diagram for cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >:

Public Types

using element_type = get_vector_element_type< Derived >
 Element type of the derived (user) class. More...
 
using raw_vector_type = vector_type_t< RawTy, N >
 The underlying raw storage vector data type. More...
 
using raw_element_type = RawTy
 The element type of the raw storage vector. More...
 

Public Member Functions

 simd_obj_impl ()=default
 Default constructor. More...
 
 simd_obj_impl (const simd_obj_impl &other)
 Copy constructor. More...
 
template<class Ty1 , typename Derived1 >
 simd_obj_impl (const simd_obj_impl< Ty1, N, Derived1, SFINAE > &other)
 Implicit conversion constructor from another simd_obj_impl object. More...
 
 simd_obj_impl (const raw_vector_type &Val)
 Implicit conversion constructor from a raw vector object. More...
 
 simd_obj_impl (Ty Val, Ty Step) noexcept
 Arithmetic progression constructor. More...
 
template<class T1 , class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
 simd_obj_impl (T1 Val) noexcept
 Broadcast constructor. More...
 
template<int N1, class = std::enable_if_t<N1 == N>>
 simd_obj_impl (const Ty(&&Arr)[N1]) noexcept
 Rvalue array-based constructor. More...
 
template<typename Flags = element_aligned_tag, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
 simd_obj_impl (const Ty *ptr, Flags={}) noexcept
 Pointer-based load constructor. More...
 
template<typename AccessorT , typename Flags = element_aligned_tag, typename = std::enable_if_t< detail::is_sycl_accessor_with<AccessorT, accessor_mode_cap::can_read, sycl::access::target::device>::value && is_simd_flag_type_v<Flags>>>
 simd_obj_impl (AccessorT acc, uint32_t offset, Flags={}) noexcept
 Accessor-based load constructor. More...
 
template<int N1>
std::enable_if_t< N1==N > copy_from (const Ty(&&Arr)[N1])
 Initializes this object from an rvalue to an array with the same number of elements. More...
 
template<typename T = simd_obj_impl, typename = sycl::detail::enable_if_t<T::length == 1>>
 operator Ty () const
 Type conversion into a scalar: <simd_obj_impl<RawTy, 1, simd<Ty,1>> to Ty. More...
 
raw_vector_type data () const
 
Derived read () const
 
Derived & write (const Derived &Val)
 Replaces the underlying data with the one taken from another object. More...
 
void merge (const Derived &Val, const simd_mask_type< N > &Mask)
 "Merges" this object's value with another object: replaces part of the underlying data with the one taken from the other object according to a mask. More...
 
void merge (const Derived &Val1, Derived Val2, const simd_mask_type< N > &Mask)
 Merges given two objects with a mask and writes resulting data into this object. More...
 
template<typename EltTy >
auto bit_cast_view () &
 Create a 1-dimensional view of this object. More...
 
template<typename EltTy , int Height, int Width>
auto bit_cast_view () &
 Create a 2-dimensional view of this object. More...
 
template<int Size, int Stride>
simd_view< Derived, region1d_t< Ty, Size, Stride > > select (uint16_t Offset=0) &
 Select elements of this object into a subregion and create a 1D view for for it. More...
 
template<int Size, int Stride>
resize_a_simd_type_t< Derived, Size > select (uint16_t Offset=0) &&
 Select and extract a subregion of this object's elements and return it as a new vector object. More...
 
Ty operator[] (int i) const
 Get value of this vector's element. More...
 
simd_view< Derived, region1d_scalar_t< Ty > > operator[] (int i)
 Return writable view of a single element. More...
 
template<int Size>
resize_a_simd_type_t< Derived, Size > iselect (const simd< uint16_t, Size > &Indices)
 Indirect select - select and extract multiple elements with given variable indices. More...
 
void iupdate (ushort Index, Ty V)
 Update single element with variable index. More...
 
template<int Size>
void iupdate (const simd< uint16_t, Size > &Indices, const resize_a_simd_type_t< Derived, Size > &Val, const simd_mask_type< Size > &Mask)
 Indirect update - update multiple elements with given variable indices. More...
 
template<int Rep>
resize_a_simd_type_t< Derived, Rep *N > replicate () const
 Replicates contents of this vector a number of times into a new vector. More...
 
template<int Rep, int W>
resize_a_simd_type_t< Derived, Rep *W > replicate_w (uint16_t Offset) const
 Shortcut to replicate_vs_w_hs with VS=0 and HS=1 to replicate a single "dense" (w/o gaps between elements) block Rep times. More...
 
template<int Rep, int VS, int W>
resize_a_simd_type_t< Derived, Rep *W > replicate_vs_w (uint16_t Offset) const
 Shortcut to replicate_vs_w_hs with HS=1 to replicate dense blocks. More...
 
template<int Rep, int VS, int W, int HS>
resize_a_simd_type_t< Derived, Rep *W > replicate_vs_w_hs (uint16_t Offset) const
 This function "replicates" a portion of this object's elements into a new object. More...
 
template<typename T1 = Ty, typename = std::enable_if_t<std::is_integral<T1>::value>>
uint16_t any () const
 See if any element is non-zero. More...
 
template<typename T1 = Ty, typename = std::enable_if_t<std::is_integral<T1>::value>>
uint16_t all () const
 See if all elements are non-zero. More...
 
template<typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE void copy_from (const Ty *addr, Flags={}) SYCL_ESIMD_FUNCTION
 Copy a contiguous block of data from memory into this simd_obj_impl object. More...
 
template<typename AccessorT , typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor< AccessorT, accessor_mode_cap::can_read, sycl::access::target::device, void > copy_from (AccessorT acc, uint32_t offset, Flags={}) SYCL_ESIMD_FUNCTION
 Copy a contiguous block of data from memory into this simd_obj_impl object. More...
 
template<typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE void copy_to (Ty *addr, Flags={}) const SYCL_ESIMD_FUNCTION
 Copy all vector elements of this object into a contiguous block in memory. More...
 
template<typename AccessorT , typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor< AccessorT, accessor_mode_cap::can_write, sycl::access::target::device, void > copy_to (AccessorT acc, uint32_t offset, Flags={}) const SYCL_ESIMD_FUNCTION
 Copy all vector elements of this object into a contiguous block in memory. More...
 
template<class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
Derived operator~ () const
 Per-element bitwise inversion, available in all subclasses, but only for integral element types (simd_mask included). More...
 
template<class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
simd_mask_type< N > operator! () const
 Unary logical negation operator, available in all subclasses, but only for integral element types (simd_mask included). More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived & operator^= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Bitwise-xor compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived & operator^= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 ^= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived & operator^= (T1 RHS)
 ^= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived & operator|= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Bitwise-or compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived & operator|= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 |= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived & operator|= (T1 RHS)
 |= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived & operator&= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Bitwise-and compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived & operator&= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 &= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived & operator&= (T1 RHS)
 &= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived & operator%= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Modulo operation compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived & operator%= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 %= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived & operator%= (T1 RHS)
 %= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> &&__ESIMD_DNS::is_simd_type_v<Derived> >>
Derived & operator<<= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Shift left compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_SHIFT_OP_FILTER >>
Derived & operator<<= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 <<= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_SHIFT_OP_FILTER >>
Derived & operator<<= (T1 RHS)
 <<= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> &&__ESIMD_DNS::is_simd_type_v<Derived> >>
Derived & operator>>= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Logical shift right compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_SHIFT_OP_FILTER >>
Derived & operator>>= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 >>= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_SHIFT_OP_FILTER >>
Derived & operator>>= (T1 RHS)
 >>= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived & operator+= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Addition operation compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived & operator+= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 += simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived & operator+= (T1 RHS)
 += scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived & operator-= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Subtraction operation compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived & operator-= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 -= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived & operator-= (T1 RHS)
 -= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived & operator*= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Multiplication operation compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived & operator*= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 *= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived & operator*= (T1 RHS)
 *= scalar version. More...
 
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived & operator/= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
 Division operation compound assignment. More...
 
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived & operator/= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
 /= simd_view version. More...
 
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived & operator/= (T1 RHS)
 /= scalar version. More...
 

Static Public Attributes

static constexpr int length = N
 The number of elements in this object. More...
 

Protected Member Functions

template<typename RTy , class ElemTy = __raw_t<typename RTy::element_type>>
ESIMD_INLINE void writeRegion (RTy Region, const vector_type_t< ElemTy, RTy::length > &Val)
 Write a simd_obj_impl-vector into a basic region of a simd_obj_impl object. More...
 
template<typename TR , typename UR , class ElemTy = __raw_t<typename TR::element_type>>
ESIMD_INLINE void writeRegion (std::pair< TR, UR > Region, const vector_type_t< ElemTy, TR::length > &Val)
 Write a simd_obj_impl-vector into a nested region of a simd_obj_impl object. More...
 
__ESIMD_DECLARE_TEST_PROXY void set (const raw_vector_type &Val)
 

Detailed Description

template<typename RawTy, int N, class Derived, class SFINAE>
class cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >

This is a base class for all ESIMD simd classes with real storage (simd, simd_mask_impl).

It wraps a clang vector as the storage for the elements. Additionally this class supports region operations that map to Intel GPU regions. The type of a region select or bit_cast_view operation is of simd_view type, which models a "window" into this object's storage and can used to read and modify it.

This class and its derivatives must follow the 'curiously recurring template' design pattern.

Template Parameters
RawTyRaw (storage) element type
NNumber of elements
Derived- A class derived from this one. Pure simd_obj_impl objects are never supposed to be constructed directly neither by user nor by ESIMD library code, instead they should always be enclosed into objects of some derived class - simd or simd_mask. This derived class is captured by this template parameter. Note that for some element types, the element type in the Derived type and this type may differ - for example, half type.
SFINAE- defaults to 'void' in the forward declarion within types.hpp, used to disable invalid specializations.

Definition at line 154 of file simd_obj_impl.hpp.

Member Typedef Documentation

◆ element_type

template<typename RawTy , int N, class Derived , class SFINAE >
using cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::element_type = get_vector_element_type<Derived>

Element type of the derived (user) class.

Definition at line 171 of file simd_obj_impl.hpp.

◆ raw_element_type

template<typename RawTy , int N, class Derived , class SFINAE >
using cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::raw_element_type = RawTy

The element type of the raw storage vector.

Definition at line 177 of file simd_obj_impl.hpp.

◆ raw_vector_type

template<typename RawTy , int N, class Derived , class SFINAE >
using cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::raw_vector_type = vector_type_t<RawTy, N>

The underlying raw storage vector data type.

Definition at line 174 of file simd_obj_impl.hpp.

Constructor & Destructor Documentation

◆ simd_obj_impl() [1/9]

template<typename RawTy , int N, class Derived , class SFINAE >
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( )
default

Default constructor.

Values of the constructed object's elements are undefined.

◆ simd_obj_impl() [2/9]

template<typename RawTy , int N, class Derived , class SFINAE >
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( const simd_obj_impl< RawTy, N, Derived, SFINAE > &  other)
inline

Copy constructor.

Parameters
otherThe other object to bitwise-copy elements from.

Definition at line 224 of file simd_obj_impl.hpp.

References __esimd_dbg_print, and cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data().

◆ simd_obj_impl() [3/9]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class Ty1 , typename Derived1 >
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( const simd_obj_impl< Ty1, N, Derived1, SFINAE > &  other)
inline

Implicit conversion constructor from another simd_obj_impl object.

Elements of the of the other object are type-converted to element_type to obtain elements of this object.

Template Parameters
Ty1Raw element type of the other object.
Derived1The actual type of the other object.
Parameters
otherThe other object.

Definition at line 236 of file simd_obj_impl.hpp.

References __esimd_dbg_print, and cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data().

◆ simd_obj_impl() [4/9]

template<typename RawTy , int N, class Derived , class SFINAE >
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( const raw_vector_type Val)
inline

Implicit conversion constructor from a raw vector object.

Parameters
Valthe raw vector to convert from.

Definition at line 243 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ simd_obj_impl() [5/9]

template<typename RawTy , int N, class Derived , class SFINAE >
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( Ty  Val,
Ty  Step 
)
inlinenoexcept

Arithmetic progression constructor.

Consecutive elements of this object are initialized with the arithmetic progression defined by the arguments. For example, simd<int, 4> x(1, 3) will initialize x to the {1, 4, 7, 10} sequence.

Parameters
ValThe start of the progression.
StepThe step of the progression.

Definition at line 254 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ simd_obj_impl() [6/9]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( T1  Val)
inlinenoexcept

Broadcast constructor.

Given value is type-converted to the element_type and resulting bit representation is broadcast to all lanes of the underlying vector.

Template Parameters
T1Type of the value.
Parameters
ValThe value to broadcast.

Definition at line 273 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ simd_obj_impl() [7/9]

template<typename RawTy , int N, class Derived , class SFINAE >
template<int N1, class = std::enable_if_t<N1 == N>>
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( const Ty(&&)  Arr[N1])
inlinenoexcept

Rvalue array-based constructor.

Used for in-place initialization like simd<int, N> x({1,0,0,1,...}).

Parameters
ArrRvalue reference to an array of size N to initialize from.

Definition at line 283 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ simd_obj_impl() [8/9]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename Flags = element_aligned_tag, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( const Ty *  ptr,
Flags  = {} 
)
inlinenoexcept

Pointer-based load constructor.

Initializes this object from values stored in memory. For example: simd<int, N> x(ptr, overaligned_tag<16>{});.

Template Parameters
FlagsSpecifies memory address alignment. Affects efficiency of the generated code.
Parameters
ptrThe memory address to read from.

Definition at line 302 of file simd_obj_impl.hpp.

◆ simd_obj_impl() [9/9]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename AccessorT , typename Flags = element_aligned_tag, typename = std::enable_if_t< detail::is_sycl_accessor_with<AccessorT, accessor_mode_cap::can_read, sycl::access::target::device>::value && is_simd_flag_type_v<Flags>>>
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl ( AccessorT  acc,
uint32_t  offset,
Flags  = {} 
)
inlinenoexcept

Accessor-based load constructor.

Initializes constructed object from values stored in memory represented by an accessor and an offset. For example: simd<int, N> x(acc, 128, overaligned_tag<16>{});.

Template Parameters
AccessorTthe type of the accessor. Auto-deduced.
FlagsSpecifies memory address alignment. Affects efficiency of the generated code. Auto-deduced from the unnamed alignment tag argument.
Parameters
accThe accessor to read from.
offset32-bit offset in bytes of the first element.

Definition at line 323 of file simd_obj_impl.hpp.

Member Function Documentation

◆ all()

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename T1 = Ty, typename = std::enable_if_t<std::is_integral<T1>::value>>
uint16_t cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::all ( ) const
inline

See if all elements are non-zero.

Returns
1 if all elements are non-zero, 0 otherwise.

Definition at line 600 of file simd_obj_impl.hpp.

◆ any()

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename T1 = Ty, typename = std::enable_if_t<std::is_integral<T1>::value>>
uint16_t cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::any ( ) const
inline

See if any element is non-zero.

Returns
1 if any element is non-zero, 0 otherwise.

Definition at line 591 of file simd_obj_impl.hpp.

◆ bit_cast_view() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename EltTy >
auto cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::bit_cast_view ( ) &
inline

Create a 1-dimensional view of this object.

Template Parameters
EltTyThe element type of the new view.
Returns
A new simd_view object which spans this entire object, but potentially with a different element type and different number of elements, if the sizes of this object's element type and the new one don't match.

Definition at line 395 of file simd_obj_impl.hpp.

◆ bit_cast_view() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename EltTy , int Height, int Width>
auto cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::bit_cast_view ( ) &
inline

Create a 2-dimensional view of this object.

sizeof(EltTy)*Height*Width must be equal to the byte size of this object.

Template Parameters
ElTyElement type of the view.
HeightHeight of the view in rows.
WidthWidth of the view in elements.
Returns
A new 2D simd_view object which spans this entire object, but potentially with a different element type and different number of elements, if the sizes of this object's element type and the new one don't match.

Definition at line 412 of file simd_obj_impl.hpp.

◆ copy_from() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename AccessorT , typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, sycl::access::target::device, void> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::copy_from ( AccessorT  acc,
uint32_t  offset,
Flags  = {} 
)

Copy a contiguous block of data from memory into this simd_obj_impl object.

The amount of memory copied equals the total size of vector elements in this object. Source memory location is represented via a global accessor and offset. None of the template parameters except documented ones can/should be specified by callers.

Template Parameters
AccessorTType of the accessor (auto-deduced).
FlagsAlignment control for the copy operation. See Alignment control. for more info.
Parameters
accaccessor to copy from.
offsetoffset to copy from (in bytes).

◆ copy_from() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::copy_from ( const Ty *  addr,
Flags  = {} 
)

Copy a contiguous block of data from memory into this simd_obj_impl object.

The amount of memory copied equals the total size of vector elements in this object. None of the template parameters except documented ones can/should be specified by callers.

Template Parameters
FlagsAlignment control for the copy operation. See Alignment control. for more info.
Parameters
addrthe memory address to copy from. Must be a pointer to the global address space, otherwise behavior is undefined.

◆ copy_from() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<int N1>
std::enable_if_t<N1 == N> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::copy_from ( const Ty(&&)  Arr[N1])
inline

Initializes this object from an rvalue to an array with the same number of elements.

Parameters
ArrRvalue reference to the array.

Definition at line 331 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ copy_to() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename AccessorT , typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, sycl::access::target::device, void> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::copy_to ( AccessorT  acc,
uint32_t  offset,
Flags  = {} 
) const

Copy all vector elements of this object into a contiguous block in memory.

Destination memory location is represented via a global accessor and offset. None of the template parameters should be be specified by callers.

Template Parameters
AccessorTType of the accessor (auto-deduced).
FlagsAlignment control for the copy operation. See Alignment control. for more info.
Parameters
accaccessor to copy from.
offsetoffset to copy from.

◆ copy_to() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename Flags = element_aligned_tag, int ChunkSize = 32, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::copy_to ( Ty *  addr,
Flags  = {} 
) const

Copy all vector elements of this object into a contiguous block in memory.

None of the template parameters should be be specified by callers.

Template Parameters
FlagsAlignment control for the copy operation. See Alignment control. for more info.
Parameters
addrthe memory address to copy to. Must be a pointer to the global address space, otherwise behavior is undefined.

◆ data()

template<typename RawTy , int N, class Derived , class SFINAE >
raw_vector_type cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data ( ) const
inline
Returns
The value of the underlying raw vector.

Definition at line 346 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

Referenced by cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::simd_obj_impl().

◆ iselect()

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Size>
resize_a_simd_type_t<Derived, Size> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::iselect ( const simd< uint16_t, Size > &  Indices)
inline

Indirect select - select and extract multiple elements with given variable indices.

Template Parameters
SizeThe number of elements to select.
Parameters
IndicesIndices of element to select.
Returns
Vector of extracted elements.

Definition at line 470 of file simd_obj_impl.hpp.

◆ iupdate() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Size>
void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::iupdate ( const simd< uint16_t, Size > &  Indices,
const resize_a_simd_type_t< Derived, Size > &  Val,
const simd_mask_type< Size > &  Mask 
)
inline

Indirect update - update multiple elements with given variable indices.

Template Parameters
SizeThe number of elements to update.
Parameters
IndicesIndices of element to update.
ValNew values.
MaskOperation mask. 1 - update, 0 - not.

Definition at line 490 of file simd_obj_impl.hpp.

◆ iupdate() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::iupdate ( ushort  Index,
Ty  V 
)
inline

Update single element with variable index.

Parameters
IndexElement index.
VNew value.

Definition at line 478 of file simd_obj_impl.hpp.

◆ merge() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::merge ( const Derived &  Val,
const simd_mask_type< N > &  Mask 
)
inline

"Merges" this object's value with another object: replaces part of the underlying data with the one taken from the other object according to a mask.

Only elements in lanes where corresponding mask's value is non-zero are replaced.

Parameters
ValThe object to take new values from.
MaskThe mask.

Definition at line 372 of file simd_obj_impl.hpp.

◆ merge() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::merge ( const Derived &  Val1,
Derived  Val2,
const simd_mask_type< N > &  Mask 
)
inline

Merges given two objects with a mask and writes resulting data into this object.

Parameters
Val1The first object, provides elements for lanes with non-zero corresponding predicates.
Val2The second object, provides elements for lanes with zero corresponding predicates.
MaskThe merge mask.

Definition at line 384 of file simd_obj_impl.hpp.

◆ operator Ty()

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename T = simd_obj_impl, typename = sycl::detail::enable_if_t<T::length == 1>>
cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator Ty ( ) const
inline

Type conversion into a scalar: <simd_obj_impl<RawTy, 1, simd<Ty,1>> to Ty.

Definition at line 340 of file simd_obj_impl.hpp.

References __esimd_dbg_print.

◆ operator!()

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
simd_mask_type<N> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator! ( ) const
inline

Unary logical negation operator, available in all subclasses, but only for integral element types (simd_mask included).

Similarly to C++, where !x returns bool, !simd returns a simd_mask, where each element is a result of comparision with zero.

Returns
A simd_mask instance where each element is a result of comparison of the original element with zero.

Definition at line 763 of file simd_obj_impl.hpp.

◆ operator%=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator%= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Modulo operation compound assignment.

Available only when elements of both this object and the argument are integral. %= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 833 of file simd_obj_impl.hpp.

◆ operator%=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator%= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

%= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 833 of file simd_obj_impl.hpp.

◆ operator%=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator%= ( T1  RHS)
inline

%= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 833 of file simd_obj_impl.hpp.

◆ operator&=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator&= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Bitwise-and compound assignment.

Available only when elements of both this object and the argument are integral. &= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 830 of file simd_obj_impl.hpp.

◆ operator&=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator&= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

&= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 830 of file simd_obj_impl.hpp.

◆ operator&=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator&= ( T1  RHS)
inline

&= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 830 of file simd_obj_impl.hpp.

◆ operator*=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator*= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Multiplication operation compound assignment.

*= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 862 of file simd_obj_impl.hpp.

◆ operator*=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator*= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

*= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 862 of file simd_obj_impl.hpp.

◆ operator*=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator*= ( T1  RHS)
inline

*= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 862 of file simd_obj_impl.hpp.

◆ operator+=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator+= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Addition operation compound assignment.

+= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 858 of file simd_obj_impl.hpp.

◆ operator+=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator+= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

+= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 858 of file simd_obj_impl.hpp.

◆ operator+=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator+= ( T1  RHS)
inline

+= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 858 of file simd_obj_impl.hpp.

◆ operator-=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator-= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Subtraction operation compound assignment.

-= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 860 of file simd_obj_impl.hpp.

◆ operator-=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator-= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

-= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 860 of file simd_obj_impl.hpp.

◆ operator-=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator-= ( T1  RHS)
inline

-= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 860 of file simd_obj_impl.hpp.

◆ operator/=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator/= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Division operation compound assignment.

/= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 864 of file simd_obj_impl.hpp.

◆ operator/=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator/= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

/= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 864 of file simd_obj_impl.hpp.

◆ operator/=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_ARITH_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator/= ( T1  RHS)
inline

/= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 864 of file simd_obj_impl.hpp.

◆ operator<<=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> &&__ESIMD_DNS::is_simd_type_v<Derived> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator<<= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Shift left compound assignment.

Available only when elements of both this object and the source are integral. Not available for simd_mask. <<= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 844 of file simd_obj_impl.hpp.

◆ operator<<=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_SHIFT_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator<<= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

<<= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 844 of file simd_obj_impl.hpp.

◆ operator<<=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_SHIFT_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator<<= ( T1  RHS)
inline

<<= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 844 of file simd_obj_impl.hpp.

◆ operator>>=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> &&__ESIMD_DNS::is_simd_type_v<Derived> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator>>= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Logical shift right compound assignment.

Available only when elements of both this object and the source are integral. Not available for simd_mask. >>= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 848 of file simd_obj_impl.hpp.

◆ operator>>=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_SHIFT_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator>>= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

>>= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 848 of file simd_obj_impl.hpp.

◆ operator>>=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_SHIFT_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator>>= ( T1  RHS)
inline

>>= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 848 of file simd_obj_impl.hpp.

◆ operator[]() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
simd_view<Derived, region1d_scalar_t<Ty> > cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator[] ( int  i)
inline

Return writable view of a single element.

Parameters
iElement index.
Returns
View of i'th element.

Definition at line 458 of file simd_obj_impl.hpp.

◆ operator[]() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
Ty cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator[] ( int  i) const
inline

Get value of this vector's element.

Parameters
iElement index.
Returns
Value of i'th element.

Definition at line 453 of file simd_obj_impl.hpp.

◆ operator^=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator^= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Bitwise-xor compound assignment.

Available only when elements of both this object and the argument are integral. ^= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 824 of file simd_obj_impl.hpp.

◆ operator^=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator^= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

^= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 824 of file simd_obj_impl.hpp.

◆ operator^=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator^= ( T1  RHS)
inline

^= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 824 of file simd_obj_impl.hpp.

◆ operator|=() [1/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class SimdT , class = std::enable_if_t<(is_simd_type_v<Derived> == is_simd_type_v<SimdT>)&& std::is_integral_v<element_type> &&std::is_integral_v<T1> >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator|= ( const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &  RHS)
inline

Bitwise-or compound assignment.

Available only when elements of both this object and the argument are integral. |= simd version.

Template Parameters
T1Element type of the argument object (auto-deduced).
SimdTThe argument object type(auto-deduced).
Parameters
RHSThe argument object.

Definition at line 827 of file simd_obj_impl.hpp.

◆ operator|=() [2/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class SimdT1 , class RegionT1 , class T1 = typename RegionT1::element_type, class = std::enable_if_t< (is_simd_type_v<Derived> == is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator|= ( const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &  RHS)
inline

|= simd_view version.


Template Parameters
SimdT1The type of the object "viewed" by the argument
(auto-deduced).
RegionT1Region type of the argument object (auto-deduced).
Parameters
RHSThe argument object.

Definition at line 827 of file simd_obj_impl.hpp.

◆ operator|=() [3/3]

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 , class = std::enable_if_t< __ESIMD_BITWISE_OP_FILTER >>
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator|= ( T1  RHS)
inline

|= scalar version.


Template Parameters
T1The type of the scalar argument (auto-deduced).
Parameters
RHSThe argument.

Definition at line 827 of file simd_obj_impl.hpp.

◆ operator~()

template<typename RawTy , int N, class Derived , class SFINAE >
template<class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
Derived cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::operator~ ( ) const
inline

Per-element bitwise inversion, available in all subclasses, but only for integral element types (simd_mask included).

Returns
Copy of this object with all elements bitwise inverted.

Definition at line 751 of file simd_obj_impl.hpp.

◆ read()

template<typename RawTy , int N, class Derived , class SFINAE >
Derived cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::read ( ) const
inline
Returns
Newly constructed (from the underlying data) object of the Derived type.

Definition at line 357 of file simd_obj_impl.hpp.

◆ replicate()

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Rep>
resize_a_simd_type_t<Derived, Rep * N> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::replicate ( ) const
inline

Replicates contents of this vector a number of times into a new vector.

Template Parameters
RepThe number of times this vector has to be replicated.
Returns
Replicated simd_obj_impl instance.

Definition at line 501 of file simd_obj_impl.hpp.

◆ replicate_vs_w()

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Rep, int VS, int W>
resize_a_simd_type_t<Derived, Rep * W> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::replicate_vs_w ( uint16_t  Offset) const
inline

Shortcut to replicate_vs_w_hs with HS=1 to replicate dense blocks.

Template Parameters
RepNumber of blocks to select for replication.
VSVertical stride - distance between first elements of consecutive blocks. If VS=0, then the same block will be replicated Rep times in the result.
WWidth - number of elements in a block.
Parameters
OffsetThe offset of the first element of the first block.
Returns
Vector of size Rep*W consisting of replicated elements of this object.

Definition at line 527 of file simd_obj_impl.hpp.

◆ replicate_vs_w_hs()

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Rep, int VS, int W, int HS>
resize_a_simd_type_t<Derived, Rep * W> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::replicate_vs_w_hs ( uint16_t  Offset) const
inline

This function "replicates" a portion of this object's elements into a new object.

The source elements to replicate are Rep number of blocks each of size W elements. Starting elements of consecutive blocks are VS elements apart and i'th block starts from ith_block_start_ind = Offset + i*VS index. Consecutive elements within a block are HS elements apart and j'th element in the block has ith_block_start_ind + j*HS index. Thus total of Rep*W elements are returned. Note that depending on VS, W and HS, blocks' elements may overlap and in this case the elements where the overlap happens may participate 2 or more times in the result.

Example 1. Source object has 32 elements, Rep is 2, VS is 17, W is 3 and HS is 4. Selected elements are depicted with their index (matching their values) instead of a dot:

simd<int, 32> Source(0/*Base*/, 1/*Step*/);
simd<int, 6> Result = Source.replicate_vs_w_hs<2,17,3,4>(1);
// |<-------------- VS=17 ------------->|
// v-------v-------v W=3
// . 1 . . . 5 . . . 9 . . . . . . . \ Rep=2
// . 18 . . . 22 . . . 26 . . . . . /
// |<- HS=4->|
// The Result is a vector of 6 source elements {1,5,9,18,22,26}.

Example 2. AOS 7x3 => SOA 3x7 conversion. Rep is 3, VS is 1, W is 7 and HS is 3.

simd<float, 21> Source = getSource();
simd<float, 21> Result = Source.replicate_vs_w_hs<3,1,7,3>(0);
// Source:
// x0 y0 z0 x1 y1 z1 x2 y2 z2 x3 y3 z3 x4 y4 z4 x5 y5 z5 x6 y6 z6
// Result:
// x0 x1 x2 x3 x4 x5 x6 y0 y1 y2 y3 y4 y5 y6 z0 z1 z2 z3 z4 z5 z6
Template Parameters
RepNumber of blocks to select for replication.
VSVertical stride - distance between first elements of consecutive blocks. If VS=0, then the same block will be replicated Rep times in the result.
WWidth - number of elements in a block.
HSHorizontal stride - distance between consecutive elements in a block.
Parameters
OffsetThe offset of the first element of the first block.
Returns
Vector of size Rep*W consisting of replicated elements of this object.

Definition at line 581 of file simd_obj_impl.hpp.

◆ replicate_w()

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Rep, int W>
resize_a_simd_type_t<Derived, Rep * W> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::replicate_w ( uint16_t  Offset) const
inline

Shortcut to replicate_vs_w_hs with VS=0 and HS=1 to replicate a single "dense" (w/o gaps between elements) block Rep times.

Template Parameters
RepThe number of times to replicate the block.
WWidth - number of elements in the block.
Parameters
OffsetOffset of the block's first element.
Returns
Vector of size Rep*W consisting of replicated elements of this object.

Definition at line 513 of file simd_obj_impl.hpp.

◆ select() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Size, int Stride>
simd_view<Derived, region1d_t<Ty, Size, Stride> > cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::select ( uint16_t  Offset = 0) &
inline

Select elements of this object into a subregion and create a 1D view for for it.

Used when this is an lvalue.

Template Parameters
SizeThe number of elements selected for the subregion.
StrideA distance in elements between two consecutive elements.
Parameters
OffsetThe starting element's offset.
Returns
A view of the subregion.

Definition at line 427 of file simd_obj_impl.hpp.

◆ select() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<int Size, int Stride>
resize_a_simd_type_t<Derived, Size> cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::select ( uint16_t  Offset = 0) &&
inline

Select and extract a subregion of this object's elements and return it as a new vector object.

Used when this is an rvalue.

Template Parameters
SizeThe number of elements selected for the subregion.
StrideA distance in elements between two consecutive elements.
Parameters
OffsetThe starting element's offset.
Returns
Extracted subregion as a new vector object.

Definition at line 442 of file simd_obj_impl.hpp.

◆ set()

template<typename RawTy , int N, class Derived , class SFINAE >
__ESIMD_DECLARE_TEST_PROXY void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::set ( const raw_vector_type Val)
inlineprotected

Definition at line 879 of file simd_obj_impl.hpp.

◆ write()

template<typename RawTy , int N, class Derived , class SFINAE >
Derived& cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::write ( const Derived &  Val)
inline

Replaces the underlying data with the one taken from another object.

Returns
This object.

Definition at line 361 of file simd_obj_impl.hpp.

◆ writeRegion() [1/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename RTy , class ElemTy = __raw_t<typename RTy::element_type>>
ESIMD_INLINE void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::writeRegion ( RTy  Region,
const vector_type_t< ElemTy, RTy::length > &  Val 
)
inlineprotected

Write a simd_obj_impl-vector into a basic region of a simd_obj_impl object.

Definition at line 608 of file simd_obj_impl.hpp.

References cl::sycl::length().

◆ writeRegion() [2/2]

template<typename RawTy , int N, class Derived , class SFINAE >
template<typename TR , typename UR , class ElemTy = __raw_t<typename TR::element_type>>
ESIMD_INLINE void cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::writeRegion ( std::pair< TR, UR >  Region,
const vector_type_t< ElemTy, TR::length > &  Val 
)
inlineprotected

Write a simd_obj_impl-vector into a nested region of a simd_obj_impl object.

Definition at line 636 of file simd_obj_impl.hpp.

References cl::sycl::length().

Member Data Documentation

◆ length

template<typename RawTy , int N, class Derived , class SFINAE >
constexpr int cl::sycl::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::length = N
staticconstexpr

The number of elements in this object.

Definition at line 180 of file simd_obj_impl.hpp.


The documentation for this class was generated from the following file:
simd
Definition: simd.hpp:1027