This class is a simd_obj_impl specialization representing a simd mask, which is basically a simd_obj_impl with fixed element type and limited set of APIs. More...

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

Inheritance diagram for sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >:

[legend]

Collaboration diagram for sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >:

[legend]

Public Types
using	raw_element_type = T
	Raw element type actually used for storage. More...

using	element_type = T
	Element type, same as raw. More...

using	raw_vector_type = typename base_type::raw_vector_type
	Underlying storage type for the entire vector. More...

Public Types inherited from sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< T, N, simd_mask_impl< T, N >, std::enable_if_t< std::is_same_v< simd_mask_elem_type, T > > >
using	element_type = get_vector_element_type< simd_mask_impl< T, N > >
	Element type of the derived (user) class. More...

using	raw_vector_type = vector_type_t< T, N >
	The underlying raw storage vector data type. More...

using	raw_element_type = T
	The element type of the raw storage vector. More...

Public Member Functions
	simd_mask_impl ()=default
	Compiler-generated default constructor. More...

	simd_mask_impl (const simd_mask_impl &other)
	Copy constructor. More...

template<class T1 , class = std::enable_if_t<std::is_integral_v<T1>>>
	simd_mask_impl (T1 Val)
	Broadcast constructor with conversion. More...

	simd_mask_impl (const raw_vector_type &Val)
	Implicit conversion constructor from a raw vector object. More...

template<int N1, class = std::enable_if_t<N1 == N>>
	simd_mask_impl (const raw_element_type(&&Arr)[N1])
	Construct from an array. To allow e.g. simd_mask<N> m({1,0,0,1,...}). More...

	simd_mask_impl (const simd< T, N > &Val)
	Implicit conversion from simd. More...

template<typename T1 >
	simd_mask_impl (const ext::oneapi::experimental::simd_mask< T1, N > &Val)
	Implicit conversion from std::experimental::simd_mask. More...

template<typename T1 , typename = std::enable_if_t<mask_size_ok_for_mem_io() && std::is_same_v<T1, element_type>>>
	simd_mask_impl (const T1 *ptr)
	Load constructor. More...

simd_mask_impl &	operator= (element_type val) noexcept
	Broadcast assignment operator to support simd_mask_impl<N> n = a > b;. More...

simd_mask_impl &	operator= (const simd_mask_impl &other) noexcept
	Copy assignment operator. More...

template<class T1 = simd_mask_impl, class = std::enable_if_t<T1::length == 1>>
	operator bool () const
	Conversion to boolean. More...

Public Member Functions inherited from sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< T, N, simd_mask_impl< T, N >, std::enable_if_t< std::is_same_v< simd_mask_elem_type, T > > >
	simd_obj_impl ()=default
	Default constructor. More...

	simd_obj_impl (const simd_obj_impl &other)
	Copy constructor. More...

	simd_obj_impl (const simd_obj_impl< Ty1, N, Derived1, std::enable_if_t< std::is_same_v< simd_mask_elem_type, T > > > &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 Base, Ty Step) noexcept
	Arithmetic progression constructor. More...

	simd_obj_impl (T1 Val) noexcept
	Broadcast constructor. More...

	simd_obj_impl (const Ty(&&Arr)[N1]) noexcept
	Rvalue array-based constructor. More...

	simd_obj_impl (const Ty *ptr, Flags) noexcept
	Pointer-based load constructor. More...

	simd_obj_impl (const Ty *ptr, PropertyListT={}) noexcept
	Pointer-based load constructor. More...

	simd_obj_impl (AccessorT acc, uint32_t offset, Flags) noexcept
	Accessor-based load constructor. More...

	simd_obj_impl (AccessorT acc, uint32_t offset, PropertyListT={}) noexcept
	Accessor-based load constructor. More...

simd_mask_impl< T, N > &	operator= (const simd_obj_impl &other) noexcept
	Copy assignment operator. More...

	operator Ty () const
	Type conversion into a scalar: `<simd_obj_impl<RawTy, 1, simd<Ty,1>>` to `Ty`. More...

raw_vector_type	data () const

raw_vector_type &	data_ref ()

simd_mask_impl< T, N >	read () const

simd_mask_impl< T, N > &	write (const simd_mask_impl< T, N > &Val)
	Replaces the underlying data with the one taken from another object. More...

void	merge (const simd_mask_impl< T, N > &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 simd_mask_impl< T, N > &Val1, simd_mask_impl< T, N > Val2, const simd_mask_type< N > &Mask)
	Merges given two objects with a mask and writes resulting data into this object. More...

auto	bit_cast_view () &
	Create a 1-dimensional view of this object. More...

auto	bit_cast_view () &
	Create a 2-dimensional view of this object. More...

simd_view< simd_mask_impl< T, N >, 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...

resize_a_simd_type_t< simd_mask_impl< T, N >, 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< simd_mask_impl< T, N >, region1d_scalar_t< Ty > >	operator[] (int i)
	Return writable view of a single element. More...

resize_a_simd_type_t< simd_mask_impl< T, N >, 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...

void	iupdate (const simd< uint16_t, Size > &Indices, const resize_a_simd_type_t< simd_mask_impl< T, N >, Size > &Val, const simd_mask_type< Size > &Mask)
	Indirect update - update multiple elements with given variable indices. More...

resize_a_simd_type_t< simd_mask_impl< T, N >, Rep *N >	replicate () const
	Replicates contents of this vector a number of times into a new vector. More...

resize_a_simd_type_t< simd_mask_impl< T, N >, 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...

resize_a_simd_type_t< simd_mask_impl< T, N >, Rep *W >	replicate_vs_w (uint16_t Offset) const
	Shortcut to `replicate_vs_w_hs` with `HS=1` to replicate dense blocks. More...

resize_a_simd_type_t< simd_mask_impl< T, N >, 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...

uint16_t	any () const
	See if any element is non-zero. More...

uint16_t	all () const
	See if all elements are non-zero. More...

ESIMD_INLINE std::enable_if_t< is_simd_flag_type_v< Flags > >	copy_from (const Ty *addr, Flags) SYCL_ESIMD_FUNCTION
	Copy a contiguous block of data from memory into this simd_obj_impl object. More...

ESIMD_INLINE std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	copy_from (const Ty *addr, PropertyListT={}) SYCL_ESIMD_FUNCTION
	Copy a contiguous block of data from memory into this simd_obj_impl object. More...

ESIMD_INLINE std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, accessor_mode_cap::can_read > &&is_simd_flag_type_v< Flags > >	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...

ESIMD_INLINE std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	copy_from (AccessorT acc, uint32_t offset, PropertyListT={}) SYCL_ESIMD_FUNCTION
	Copy a contiguous block of data from memory into this simd_obj_impl object. More...

ESIMD_INLINE std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, accessor_mode_cap::can_read > &&is_simd_flag_type_v< Flags >, 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...

ESIMD_INLINE std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, void >	copy_from (AccessorT acc, uint32_t offset, PropertyListT={}) SYCL_ESIMD_FUNCTION
	Copy a contiguous block of data from memory into this simd_obj_impl object. More...

ESIMD_INLINE std::enable_if_t< is_simd_flag_type_v< Flags > >	copy_to (Ty *addr, Flags) const SYCL_ESIMD_FUNCTION
	Copy all vector elements of this object into a contiguous block in memory. More...

ESIMD_INLINE std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	copy_to (Ty *addr, PropertyListT={}) const SYCL_ESIMD_FUNCTION
	Copy all vector elements of this object into a contiguous block in memory. More...

ESIMD_INLINE std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, accessor_mode_cap::can_write > &&is_simd_flag_type_v< Flags > >	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...

ESIMD_INLINE std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	copy_to (AccessorT acc, uint32_t offset, PropertyListT={}) const SYCL_ESIMD_FUNCTION
	Copy all vector elements of this object into a contiguous block in memory. More...

ESIMD_INLINE std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, accessor_mode_cap::can_write > &&is_simd_flag_type_v< Flags >, 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...

ESIMD_INLINE std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, void >	copy_to (AccessorT acc, uint32_t offset, PropertyListT={}) const SYCL_ESIMD_FUNCTION
	Copy all vector elements of this object into a contiguous block in memory. More...

simd_mask_impl< T, N >	operator~ () const
	Per-element bitwise inversion, available in all subclasses, but only for integral element types (`simd_mask` included). More...

simd_mask_type< N >	operator! () const
	Unary logical negation operator, available in all subclasses, but only for integral element types (`simd_mask` included). More...

simd_mask_impl< T, N > &	operator^= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Bitwise-`xor` compound assignment. More...

simd_mask_impl< T, N > &	operator^= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`^=` simd_view version. More...

simd_mask_impl< T, N > &	operator^= (T1 RHS)
	`^=` scalar version. More...

simd_mask_impl< T, N > &	operator\|= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Bitwise-`or` compound assignment. More...

simd_mask_impl< T, N > &	operator\|= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	\|= simd_view version. More...

simd_mask_impl< T, N > &	operator\|= (T1 RHS)
	\|= scalar version. More...

simd_mask_impl< T, N > &	operator&= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Bitwise-`and` compound assignment. More...

simd_mask_impl< T, N > &	operator&= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`&=` simd_view version. More...

simd_mask_impl< T, N > &	operator&= (T1 RHS)
	`&=` scalar version. More...

simd_mask_impl< T, N > &	operator%= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Modulo operation compound assignment. More...

simd_mask_impl< T, N > &	operator%= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`%=` simd_view version. More...

simd_mask_impl< T, N > &	operator%= (T1 RHS)
	`%=` scalar version. More...

simd_mask_impl< T, N > &	operator<<= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Shift left compound assignment. More...

simd_mask_impl< T, N > &	operator<<= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`<<=` simd_view version. More...

simd_mask_impl< T, N > &	operator<<= (T1 RHS)
	`<<=` scalar version. More...

simd_mask_impl< T, N > &	operator>>= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Logical shift right compound assignment. More...

simd_mask_impl< T, N > &	operator>>= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`>>=` simd_view version. More...

simd_mask_impl< T, N > &	operator>>= (T1 RHS)
	`>>=` scalar version. More...

simd_mask_impl< T, N > &	operator+= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Addition operation compound assignment. More...

simd_mask_impl< T, N > &	operator+= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`+=` simd_view version. More...

simd_mask_impl< T, N > &	operator+= (T1 RHS)
	`+=` scalar version. More...

simd_mask_impl< T, N > &	operator-= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Subtraction operation compound assignment. More...

simd_mask_impl< T, N > &	operator-= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`-=` simd_view version. More...

simd_mask_impl< T, N > &	operator-= (T1 RHS)
	`-=` scalar version. More...

simd_mask_impl< T, N > &	operator*= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Multiplication operation compound assignment. More...

simd_mask_impl< T, N > &	operator*= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`*=` simd_view version. More...

simd_mask_impl< T, N > &	operator*= (T1 RHS)
	`*=` scalar version. More...

simd_mask_impl< T, N > &	operator/= (const __ESIMD_DNS::simd_obj_impl< T1, N, SimdT > &RHS)
	Division operation compound assignment. More...

simd_mask_impl< T, N > &	operator/= (const sycl::ext::intel::esimd ::simd_view< SimdT1, RegionT1 > &RHS)
	`/=` simd_view version. More...

simd_mask_impl< T, N > &	operator/= (T1 RHS)
	`/=` scalar version. More...

Additional Inherited Members
Static Public Attributes inherited from sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< T, N, simd_mask_impl< T, N >, std::enable_if_t< std::is_same_v< simd_mask_elem_type, T > > >
static constexpr int	length
	The number of elements in this object. More...

Protected Member Functions inherited from sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< T, N, simd_mask_impl< T, N >, std::enable_if_t< std::is_same_v< simd_mask_elem_type, T > > >
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...

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...

void	set (const raw_vector_type &Val)

Detailed Description

template<typename T, int N>
class sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >

This class is a simd_obj_impl specialization representing a simd mask, which is basically a simd_obj_impl with fixed element type and limited set of APIs.

E.g. arithmetic operations (+, -, etc.) are not defined for masks, but bit operations like ^, & are. Masks are used in many ESIMD APIs to enable/disable specific lanes - for example, to disable certain lanes when writing to a memory via the scatter operation. Each mask element may thus be interpreted in one of two ways - by ESIMD ops - "lane enaled" or "lane disabled". They can be produced in a number of ways:

using per-element simd object comparison operations, in which case the result is true / false per lane, and true means "enabled" in the resulting mask lane, false - "disabled".
reading from memory, for example:
using simd_mask_elem_t = typename simd_mask<1>::element_type;

simd_mask_elem_t *arr;

...

simd_mask<8> m(arr + 8*i);

In this and all other cases below lanes with non-zero value are treated as "enabled", lanes with zero - as "disabled".
constructing from an array or using broadcast constructor:
simd_mask<8> m1({1,0,1,0,1,0,1,0});

simd_mask<8> m2(1); // all "enabled"

simd_mask
Definition: simd.hpp:1651
constructing from a simd object. User code should use simd_mask<1>::element_type<code> when the mask element type needs to be used - for example, to declare a pointer to memory where mask elements can be written to/read from. Yet it must not assume it to be of any specific type (which is unsigned 16-bit integer in fact). Template Parameters T Fixed element type, must be simd_mask_elem_type. N Number of elements (per-lane predicates) in the mask.

Definition at line 57 of file simd_mask_impl.hpp.

Member Typedef Documentation

◆ element_type

template<typename T , int N>

using sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::element_type = T

Element type, same as raw.

Definition at line 69 of file simd_mask_impl.hpp.

◆ raw_element_type

template<typename T , int N>

using sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::raw_element_type = T

Raw element type actually used for storage.

Definition at line 67 of file simd_mask_impl.hpp.

◆ raw_vector_type

template<typename T , int N>

using sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::raw_vector_type = typename base_type::raw_vector_type

Underlying storage type for the entire vector.

Definition at line 71 of file simd_mask_impl.hpp.

Constructor & Destructor Documentation

◆ simd_mask_impl() [1/8]

template<typename T , int N>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( )

default

Compiler-generated default constructor.

◆ simd_mask_impl() [2/8]

template<typename T , int N>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const simd_mask_impl< T, N > & other )

inline

Copy constructor.

Definition at line 79 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [3/8]

template<typename T , int N>

template<class T1 , class = std::enable_if_t<std::is_integral_v<T1>>>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( T1 Val )

inline

Broadcast constructor with conversion.

See also: simd_obj_impl::simd_obj_impl(T)

Definition at line 84 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [4/8]

template<typename T , int N>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const raw_vector_type & Val )

inline

Implicit conversion constructor from a raw vector object.

Definition at line 88 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [5/8]

template<typename T , int N>

template<int N1, class = std::enable_if_t<N1 == N>>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const raw_element_type(&&) Arr[N1] )

inline

Construct from an array. To allow e.g. simd_mask<N> m({1,0,0,1,...}).

Definition at line 92 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [6/8]

template<typename T , int N>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const simd< T, N > & Val )

inline

Implicit conversion from simd.

Definition at line 97 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [7/8]

template<typename T , int N>

template<typename T1 >

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const ext::oneapi::experimental::simd_mask< T1, N > & Val )

inline

Implicit conversion from std::experimental::simd_mask.

Definition at line 101 of file simd_mask_impl.hpp.

◆ simd_mask_impl() [8/8]

template<typename T , int N>

template<typename T1 , typename = std::enable_if_t<mask_size_ok_for_mem_io() && std::is_same_v<T1, element_type>>>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::simd_mask_impl ( const T1 * ptr )

inlineexplicit

Load constructor.

Definition at line 125 of file simd_mask_impl.hpp.

Member Function Documentation

◆ operator bool()

template<typename T , int N>

template<class T1 = simd_mask_impl, class = std::enable_if_t<T1::length == 1>>

sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::operator bool ( ) const

inline

Conversion to boolean.

Available only when the number of elements is 1.

Returns: true if the element is non-zero, false otherwise.

Definition at line 144 of file simd_mask_impl.hpp.

◆ operator=() [1/2]

template<typename T , int N>

simd_mask_impl& sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::operator= ( const simd_mask_impl< T, N > & other )

inlinenoexcept

Copy assignment operator.

Definition at line 136 of file simd_mask_impl.hpp.

References sycl::_V1::ext::oneapi::experimental::operator=().

◆ operator=() [2/2]

template<typename T , int N>

simd_mask_impl& sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >::operator= ( element_type val )

inlinenoexcept

Broadcast assignment operator to support simd_mask_impl<N> n = a > b;.

Definition at line 130 of file simd_mask_impl.hpp.

The documentation for this class was generated from the following file:

include/sycl/ext/intel/esimd/detail/simd_mask_impl.hpp

T	Fixed element type, must be simd_mask_elem_type.
N	Number of elements (per-lane predicates) in the mask.

Public Types

Public Member Functions

Additional Inherited Members

Detailed Description

template<typename T, int N> class sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >

Member Typedef Documentation

◆ element_type

◆ raw_element_type

◆ raw_vector_type

Constructor & Destructor Documentation

◆ simd_mask_impl() [1/8]

◆ simd_mask_impl() [2/8]

◆ simd_mask_impl() [3/8]

◆ simd_mask_impl() [4/8]

◆ simd_mask_impl() [5/8]

◆ simd_mask_impl() [6/8]

◆ simd_mask_impl() [7/8]

◆ simd_mask_impl() [8/8]

Member Function Documentation

◆ operator bool()

◆ operator=() [1/2]

◆ operator=() [2/2]

template<typename T, int N>
class sycl::_V1::ext::intel::esimd::detail::simd_mask_impl< T, N >