Memory access API.

ESIMD APIs to access memory via accessors, USM pointers, perform per-element atomic operations. More...

Collaboration diagram for Memory access API.:

Modules
	Atomic memory access.
	Memory access functions which perform per-lane atomic update using given operation.
	Shared local memory access functions.
	Block load/prefetch/store functions.
	LSC-specific memory access APIs.
	This group combines types and functions specific to LSC, which is available in Intel GPUs starting from PVC and ACM.
	Named barrier APIs.
	HW thread .
	Raw send APIs.
	Implements the send instruction to send messages to variaous components of the Intel(R) processor graphics, as defined in the documentation at https://www.intel.com/content/www/us/en/docs/graphics-for-linux/developer-reference/1-0/hardware-specs.html.

Enumerations
enum	sycl::_V1::ext::intel::esimd::fence_mask : uint8_t {   sycl::_V1::ext::intel::esimd::global_coherent_fence = 0x1 , sycl::_V1::ext::intel::esimd::l2_flush_instructions = 0x2 , sycl::_V1::ext::intel::esimd::__SYCL_DEPRECATED , sycl::_V1::ext::intel::esimd::l2_flush_texture_data = 0x4 ,   sycl::_V1::ext::intel::esimd::__SYCL_DEPRECATED , sycl::_V1::ext::intel::esimd::l2_flush_constant_data = 0x8 , sycl::_V1::ext::intel::esimd::__SYCL_DEPRECATED , sycl::_V1::ext::intel::esimd::l2_flush_rw_data = 0x10 ,   sycl::_V1::ext::intel::esimd::__SYCL_DEPRECATED , sycl::_V1::ext::intel::esimd::local_barrier = 0x20 , sycl::_V1::ext::intel::esimd::l1_flush_ro_data = 0x40 , sycl::_V1::ext::intel::esimd::__SYCL_DEPRECATED }
	Represetns a bit mask to control behavior of esimd::fence. More...

Functions
template<typename AccessorTy >
__ESIMD_API SurfaceIndex	sycl::_V1::ext::intel::esimd::get_surface_index (AccessorTy acc)
	Get surface index corresponding to a SYCL accessor. More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-1) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-4) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (usm-ga-5) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-ga-6) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-7) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-8) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (const T *p, OffsetSimdViewT byte_offsets, PropertyListT props={})
	simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-ga-9) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< Tx, N > >	sycl::_V1::ext::intel::esimd::gather (const Tx *p, Toffset offset, simd_mask< N > mask=1)
	A variation of gather API with offsets represented as scalar. More...
template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-1) More...
template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	Writes ("scatters") elements of the input vector to different memory locations. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Writes ("scatters") elements of the input vector to different memory locations. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (usm-sc-4) More...
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	sycl::_V1::ext::intel::esimd::scatter (Tx *p, Toffset offset, simd< Tx, N > vals, simd_mask< N > mask=1)
	A variation of scatter API with offsets represented as scalar. More...
template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags > >	sycl::_V1::ext::intel::esimd::block_store (Tx *addr, simd< Tx, N > vals, Flags)
	Stores elements of the vector vals to a contiguous block of memory at the given address addr. More...
template<typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > byte_offsets, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
	Accessor-based gather. More...
template<typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, detail::DeviceAccessorOffsetT glob_offset)
	Loads and broadcasts the element located at acc and byte offset glob_offset to a vector and returns it as a simd object. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-1) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-4) This function is identical to (acc-ga-1) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (acc-ga-5) This function is identical to (acc-ga-2) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-ga-6) This function is identical to (acc-ga-3) except that vector size is fixed to 1. More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Accessor-based scatter. More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-2) More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-3) More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-4) More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) &&detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< detail::DeviceAccessorOffsetT, N > offsets, simd< T, N > vals, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
	Writes elements of a simd object into an accessor at given offsets. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) &&detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > >	sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, detail::DeviceAccessorOffsetT glob_offset, simd< T, N > vals, simd_mask< N > mask=1)
template<typename T , typename AccessorTy >
__ESIMD_API T	sycl::_V1::ext::intel::esimd::scalar_load (AccessorTy acc, detail::DeviceAccessorOffsetT offset)
	Load a scalar value from an accessor. More...
template<typename T , typename AccessorTy >
__ESIMD_API void	sycl::_V1::ext::intel::esimd::scalar_store (AccessorTy acc, detail::DeviceAccessorOffsetT offset, T val)
	Store a scalar value into an accessor. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API simd< T, N *get_num_channels_enabled(RGBAMask)>	sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, simd< Toffset, N > offsets, simd_mask< N > mask=1)
	Gather and transpose pixels from given memory locations defined by the base pointer p and offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, OffsetSimdViewT offsets, simd_mask< N > mask=1)
	A variation of gather_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, Toffset offset, simd_mask< N > mask=1)
	A variation of gather_rgba API with offsets represented as scalar. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API void	sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd< Toffset, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	Transpose and scatter pixels to given memory locations defined by the base pointer p and offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > >	sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, OffsetSimdViewT offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	A variation of scatter_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, Toffset offset, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	A variation of scatter_rgba API with offsets represented as scalar. More...
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 >	sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<((N==8||N==16||N==32) &&sizeof(T)==4 &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >), simd< T, N *get_num_channels_enabled(RGBAMask)> >	sycl::_V1::ext::intel::esimd::gather_rgba (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > offsets, detail::DeviceAccessorOffsetT global_offset=0, simd_mask< N > mask=1)
	Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	sycl::_V1::ext::intel::esimd::scatter_rgba (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, detail::DeviceAccessorOffsetT global_offset=0, simd_mask< N > mask=1)
	Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
template<uint8_t cntl>
__ESIMD_API void	sycl::_V1::ext::intel::esimd::fence ()
	esimd::fence sets the memory read/write order. More...
__ESIMD_API void	sycl::_V1::ext::intel::esimd::fence (fence_mask cntl)
template<memory_kind Kind = memory_kind::global, fence_flush_op FenceOp = fence_flush_op::none, fence_scope Scope = fence_scope::group>
__ESIMD_API void	sycl::_V1::ext::intel::esimd::fence ()
	Memory fence. More...
__ESIMD_API void	sycl::_V1::ext::intel::esimd::barrier ()
	Generic work-group barrier. More...
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API simd< T, m *N >	sycl::_V1::ext::intel::esimd::media_block_load (AccessorTy acc, unsigned x, unsigned y)
	Media block load. More...
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API void	sycl::_V1::ext::intel::esimd::media_block_store (AccessorTy acc, unsigned x, unsigned y, simd< T, m *N > vals)
	Media block store. More...
template<typename T , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read > &&is_simd_flag_type_v< Flags >, simd< T, N > >	sycl::_V1::ext::intel::esimd::block_load (AccessorTy acc, uint32_t byte_offset, Flags flags)
	Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte_offset, then returns the loaded data as a simd object. More...
template<typename T , int N, typename AccessorT , typename Flags >
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&is_simd_flag_type_v< Flags > >	sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, uint32_t offset, simd< T, N > vals, Flags flags)
	Variant of block_store that uses local accessor as a parameter. More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	Variant of gather that uses local accessor as a parameter template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-1) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-4) This function is identical to (lacc-ga-1) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (lacc-ga-5) This function is identical to (lacc-ga-2) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (lacc-ga-6) This function is identical to (lacc-ga-3) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read >, simd< T, N > >	sycl::_V1::ext::intel::esimd::gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset, simd_mask< N > mask=1)
	Variant of gather that uses local accessor as a parameter. More...
template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Variant of scatter that uses local accessor as a parameter template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-1) More...
template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-2) More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-3) More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-4) More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-1) void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (usm-pf-3) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-pf-4) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-5) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-pf-6) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetT byte_offset, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-7) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetT byte_offset, PropertyListT props={})
	template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}); // (usm-pf-8) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}); //(usm-pf-9) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (const T *p, PropertyListT props={})
	template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, PropertyListT props = {}); // (usm-pf-10) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-1) void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2) More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (acc-pf-3) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (acc-pf-4) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-5) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (acc-pf-6) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, OffsetT byte_offset, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (acc-pf-7) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, OffsetT byte_offset, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, PropertyListT props = {}); // (acc-pf-8) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd_mask<1> mask, PropertyListT props = {}); //(acc-pf-9) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch (AccessorT acc, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, PropertyListT props = {}); // (acc-pf-10) Supported platforms: DG2, PVC only. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	sycl::_V1::ext::intel::esimd::load_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props={})
	template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = empty_properties_t> simd<T, N> load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block load. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::prefetch_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props={})
	template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>(), typename PropertyListT = empty_properties_t> void prefetch_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block prefetch. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	sycl::_V1::ext::intel::esimd::store_2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, simd< T, N > Vals, PropertyListT props={})
	2D USM pointer block store. More...
template<split_barrier_action flag>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::split_barrier ()
	Generic work-group split barrier. More...
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::split_barrier (split_barrier_action flag)

Detailed Description

ESIMD APIs to access memory via accessors, USM pointers, perform per-element atomic operations.

Enumeration Type Documentation

fence_mask

enum sycl::_V1::ext::intel::esimd::fence_mask : uint8_t

Represetns a bit mask to control behavior of esimd::fence.

Enum elements define semantics of the bits in the mask.

Enumerator
global_coherent_fence	“Commit enable” - wait for fence to complete before continuing.
l2_flush_instructions	Flush the instruction cache.
__SYCL_DEPRECATED	Creates a software (compiler) barrier, which does not generate any instruction and only prevents instruction scheduler from reordering instructions across this barrier at compile time.
l2_flush_texture_data	Flush sampler (texture) cache.
__SYCL_DEPRECATED	Creates a software (compiler) barrier, which does not generate any instruction and only prevents instruction scheduler from reordering instructions across this barrier at compile time.
l2_flush_constant_data	Flush constant cache.
__SYCL_DEPRECATED	Creates a software (compiler) barrier, which does not generate any instruction and only prevents instruction scheduler from reordering instructions across this barrier at compile time.
l2_flush_rw_data	Flush constant cache.
__SYCL_DEPRECATED	Creates a software (compiler) barrier, which does not generate any instruction and only prevents instruction scheduler from reordering instructions across this barrier at compile time.
local_barrier	Issue SLM memory barrier only. If not set, the memory barrier is global.
l1_flush_ro_data	Flush L1 read - only data cache.
__SYCL_DEPRECATED	Creates a software (compiler) barrier, which does not generate any instruction and only prevents instruction scheduler from reordering instructions across this barrier at compile time.

Definition at line 7886 of file memory.hpp.

Function Documentation

barrier()

__ESIMD_API void sycl::_V1::ext::intel::esimd::barrier

(

)

Generic work-group barrier.

Performs barrier synchronization for all threads within the same thread group. The barrier instruction causes the executing thread to wait until all threads in the same thread group have executed the barrier instruction. Memory ordering is also guaranteed by this instruction. The behavior is undefined if this instruction is executed in divergent control flow.

Definition at line 7951 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::global_coherent_fence, and sycl::_V1::ext::intel::esimd::local_barrier.

Referenced by sycl::_V1::ext::intel::experimental::esimd::named_barrier_signal().

block_load()

template<typename T , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read> && is_simd_flag_type_v<Flags>, simd<T, N> > sycl::_V1::ext::intel::esimd::block_load	(	AccessorTy	acc,
		uint32_t	byte_offset,
		Flags	flags
	)

Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte_offset, then returns the loaded data as a simd object.

The generated code depends on the combination {T, N, Flags}. Providing flags specifying the alignment of 16-bytes or more produces more efficient code. If the alignment is smaller than 16-bytes, then less efficient gather is generated. If the loaded vector is too long for 1 flat-load GPU instruction, then a series of flat-loads and/or gathers may be generated.

Template Parameters

T	Element type.
N	Number of elements to load.
AccessorTy	Accessor type (auto-deduced).
Flags	The alignment specifier type tag.

Parameters

acc	The local accessor.
byte_offset	The offset to load from in bytes.
Flags	Specifies the alignment.

Returns

A vector of loaded elements.

Definition at line 8068 of file memory.hpp.

block_store() [1/2]

template<typename T , int N, typename AccessorT , typename Flags >

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write> && is_simd_flag_type_v<Flags> > sycl::_V1::ext::intel::esimd::block_store	(	AccessorT	acc,
		uint32_t	offset,
		simd< T, N >	vals,
		Flags	flags
	)

Variant of block_store that uses local accessor as a parameter.

Stores elements of the vector vals to a contiguous block of SLM memory represented by the given local accessor and the byte-offset offset. The generated code depends on the combination {T, N, Flags}. Providing flags specifying the alignment of 16-bytes or more produces more efficient code. If the alignment is smaller than 16-bytes, then less efficient scatter is generated. If the stored vector is too long for 1 flat-store GPU instruction, then a series of flat-store and/or scatters may be generated.

Template Parameters

T	Element type.
N	Number of elements to store.
AccessorT	Accessor type (auto-deduced).

Parameters

acc	The local accessor to store to.
offset	The byte-offset to store at.
vals	The vector to store.
Flags	Specifies the alignment.

Definition at line 8095 of file memory.hpp.

block_store() [2/2]

template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>

__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags> > sycl::_V1::ext::intel::esimd::block_store	(	Tx *	addr,
		simd< Tx, N >	vals,
		Flags
	)

Stores elements of the vector vals to a contiguous block of memory at the given address addr.

The generated code depends on the combination {T, N, Flags}. Providing flags specifying the alignment of 16-bytes or more produces more efficient code. If the alignment is smaller than 16-bytes, then less efficient scatter is generated. If the stored vector is too long for 1 flat-store GPU instruction, then a series of flat-store and/or scatters may be generated.

Template Parameters

Tx	Element type.
N	Number of elements to store.
Flags	The alignment specifier type tag.

Parameters

addr	The memory address to store at.
vals	The vector to store.
Flags	Specifies the alignment.

Definition at line 1277 of file memory.hpp.

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

fence() [1/3]

template<uint8_t cntl>

__ESIMD_API void sycl::_V1::ext::intel::esimd::fence

(

)

esimd::fence sets the memory read/write order.

Template Parameters

cntl	A bitmask composed from fence_mask bits.

Definition at line 7918 of file memory.hpp.

fence() [2/3]

template<memory_kind Kind = memory_kind::global, fence_flush_op FenceOp = fence_flush_op::none, fence_scope Scope = fence_scope::group>

__ESIMD_API void sycl::_V1::ext::intel::esimd::fence

(

)

Memory fence.

Supported platforms: DG2, PVC

Template Parameters

Kind	is the memory kind.
FenceOp	is the fence cache flush operation to apply after fence.
Scope	is the fence operation scope.

Definition at line 7932 of file memory.hpp.

References simd_mask< _Tp, _Abi >::data(), sycl::_V1::ext::intel::esimd::group, sycl::_V1::ext::intel::esimd::local, and sycl::_V1::ext::intel::esimd::none.

fence() [3/3]

__ESIMD_API void sycl::_V1::ext::intel::esimd::fence

(

fence_mask

cntl

)

Definition at line 7921 of file memory.hpp.

gather() [1/28]

template<typename T , int N, typename AccessorT >

__ESIMD_API std::enable_if_t<detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		detail::DeviceAccessorOffsetT	glob_offset
	)

Loads and broadcasts the element located at acc and byte offset glob_offset to a vector and returns it as a simd object.

Supported platforms: DG2/PVC if sizeof(T) > 4 or the number of elements to load is not equal to 1, 8, 16, 32. Otherwise, it is supported on ALL platforms.

Template Parameters

T	Element type.
N	The number of vector elements.
AccessorT	The accessor type.

Parameters

acc	The accessor to gather from.
glob_offset	Offset in bytes added to each individual element's offset to compute actual memory access offset for that element.

Definition at line 3259 of file memory.hpp.

gather() [2/28]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (acc-ga-9) This function is identical to (acc-ga-3) except that the byte_offsets is represented as simd_view.

Definition at line 3602 of file memory.hpp.

gather() [3/28]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-8) This function is identical to (acc-ga-2) except that the byte_offsets is represented as simd_view.

Definition at line 3580 of file memory.hpp.

gather() [4/28]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-7) This function is identical to (acc-ga-1) except that the byte_offsets is represented as simd_view.

Definition at line 3557 of file memory.hpp.

gather() [5/28]

template<typename T , int N, typename AccessorT >

__ESIMD_API std::enable_if_t<detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< detail::DeviceAccessorOffsetT, N >	byte_offsets,
		detail::DeviceAccessorOffsetT	glob_offset,
		simd_mask< N >	mask = 1
	)

Accessor-based gather.

Collects elements from memory referenced by the accessor acc, byte offsets byte_offsets and common offset \glob_offset, then returns the loaded elements as a single simd object.

Supported platforms: DG2/PVC if sizeof(T) > 4 or the number of elements to load is not equal to 1, 8, 16, 32. Otherwise, it is supported on ALL platforms.

Template Parameters

T	Element type.
N	The number of vector elements.
AccessorT	The accessor type.

Parameters

acc	The accessor to gather from.
byte_offsets	Per-element offsets in bytes.
glob_offset	Offset in bytes added to each individual element's offset to compute actual memory access offset for that element.
mask	Memory access mask. Elements with zero corresponding mask's predicate are not accessed, their values in the resulting vector are undefined.

Definition at line 3221 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::default_size, sycl::_V1::ext::intel::esimd::detail::gather_impl(), and sycl::_V1::ext::intel::esimd::detail::isPowerOf2().

gather() [6/28]

template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && std::is_same_v<MaskT, simd_mask<N>> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N >	byte_offsets,
		MaskT	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (acc-ga-5) This function is identical to (acc-ga-2) except that vector size is fixed to 1.

This variant is added for convenience and let user omit the template arguments and call the function as 'gather(acc, byte_offsets, mask);'.

Definition at line 3515 of file memory.hpp.

gather() [7/28]

template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && std::is_same_v<MaskT, simd_mask<N>> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N >	byte_offsets,
		MaskT	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-4) This function is identical to (acc-ga-1) except that vector size is fixed to 1.

This variant is added for convenience and let user omit the template arguments and call the function as 'gather(acc, byte_offsets, mask, pass_thru);'.

Definition at line 3489 of file memory.hpp.

gather() [8/28]

template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-ga-6) This function is identical to (acc-ga-3) except that vector size is fixed to 1.

This variant is added for convenience and let user omit the template arguments and call the function as 'gather(acc, byte_offsets);'.

Definition at line 3535 of file memory.hpp.

gather() [9/28]

template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) Supported platforms: DG2, PVC in most cases.

DG2/PVC is not required if VS == 1 and no L1/L2 cache hints used and sizeof(T) <= 4 and N = {1,8,16,32}

Loads ("gathers") elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, and returns the loaded elements.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 3460 of file memory.hpp.

gather() [10/28]

template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) Supported platforms: DG2, PVC in most cases.

DG2/PVC is not required if VS == 1 and no L1/L2 cache hints used and sizeof(T) <= 4 and N = {1,8,16,32}

Loads ("gathers") elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (acc + byte_offsets[i]) is skipped and the corresponding i-th element of the returned vector is undefined.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 3406 of file memory.hpp.

gather() [11/28]

template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-1) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3)

The next 3 functions are similar to (acc-ga-1,2,3), but they don't have the template parameter 'VS'. These functions are added for convenience and to make it possible for user to omit the template parameters T and N, e.g. 'auto res = gather(acc, byte_offsets); template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-4) simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (acc-ga-5) simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-ga-6)

The next 3 functions are similar to (acc-ga-1,2,3), but accept the byte_offsets as a simd_view argument: template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-7) simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-8) simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (acc-ga-9) template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-1) Supported platforms: DG2, PVC only - Temporary restriction for the variant with pass_thru operand. The only exception: DG2/PVC is not required if stateless memory mode is enforced via -fsycl-esimd-force-stateless-mem and VS == 1 and no L1/L2 cache hints passed and the __ESIMD_GATHER_SCATTER_LLVM_IR macro is used.

Loads ("gathers") elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (acc + byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
pass_thru	The vector pass through values.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 3357 of file memory.hpp.

gather() [12/28]

template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && std::is_same_v<MaskT, simd_mask<N>> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N >	byte_offsets,
		MaskT	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (lacc-ga-5) This function is identical to (lacc-ga-2) except that vector size is fixed to 1.

This variant is added for convenience and let user omit the template arguments and call the function as 'gather(acc, byte_offsets, mask);'.

Definition at line 8317 of file memory.hpp.

gather() [13/28]

template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && std::is_same_v<MaskT, simd_mask<N>> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N >	byte_offsets,
		MaskT	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-4) This function is identical to (lacc-ga-1) except that vector size is fixed to 1.

This variant is added for convenience and lets the user omit the template arguments and call the function as 'gather(acc, byte_offsets, mask, pass_thru);'.

Definition at line 8290 of file memory.hpp.

gather() [14/28]

template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (lacc-ga-6) This function is identical to (lacc-ga-3) except that vector size is fixed to 1.

This variant is added for convenience and let user omit the template arguments and call the function as 'gather(acc, byte_offsets);'.

Definition at line 8338 of file memory.hpp.

gather() [15/28]

template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N/VS >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) Supported platforms: DG2, PVC in most cases.

DG2/PVC is not required if VS == 1 and the __ESIMD_GATHER_SCATTER_LLVM_IR macro is used or sizeof(T) <= 4 and N = {1,2,4,8,16,32}

Loads ("gathers") elements of the type 'T' from memory locations addressed by the local accessor acc and byte offsets byte_offsets, and returns the loaded elements.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 8262 of file memory.hpp.

gather() [16/28]

template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) Supported platforms: DG2, PVC in most cases.

DG2/PVC is not required if VS == 1 and the __ESIMD_GATHER_SCATTER_LLVM_IR macro is used or sizeof(T) <= 4 and N = {1,2,4,8,16,32}

Loads ("gathers") elements of the type 'T' from memory locations addressed by the local accessor acc and byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (acc + byte_offsets[i]) is skipped and the corresponding i-th element of the returned vector is undefined.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' property is used.

Returns

A vector of elements read.

Definition at line 8224 of file memory.hpp.

gather() [17/28]

template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		simd< uint32_t, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

Variant of gather that uses local accessor as a parameter template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-1) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3)

The next 3 functions are similar to (lacc-ga-1,2,3), but they don't have the template parameter 'VS'. These functions are added for convenience and to make it possible for the user to omit the template parameters T and N, e.g. 'auto res = gather(acc, byte_offsets); template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-4) simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});//(lacc-ga-5) simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (lacc-ga-6)

The next 3 functions are similar to (lacc-ga-1,2,3), but accept the byte_offsets as a simd_view argument: template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-7) simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-8) simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (lacc-ga-9) template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-1) Supported platforms: DG2, PVC only - Temporary restriction for the variant with pass_thru operand. The only exception: DG2/PVC is not required if the __ESIMD_GATHER_SCATTER_LLVM_IR macro is used.

Loads ("gathers") elements of the type 'T' from memory locations addressed by the local accessor acc and byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (acc + byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of 32-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
pass_thru	The vector pass through values.
props	The optional compile-time properties. Only 'alignment' property is used.

Returns

A vector of elements read.

Definition at line 8179 of file memory.hpp.

gather() [18/28]

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorTy	acc,
		simd< uint32_t, N >	offsets,
		uint32_t	glob_offset,
		simd_mask< N >	mask = 1
	)

Variant of gather that uses local accessor as a parameter.

Collects elements located at given offsets in an accessor and returns them as a single simd object. An element can be a 1, 2 or 4-byte value.

Template Parameters

T	Element type; can only be a 1,2,4-byte integer, sycl::half or float.
N	The number of vector elements. Can be 1, 8, 16 or 32.
AccessorTy	The accessor type.

Parameters

acc	The accessor to gather from.
offsets	Per-element offsets in bytes.
glob_offset	Offset in bytes added to each individual element's offset to compute actual memory access offset for that element.
mask	Memory access mask. Elements with zero corresponding mask's predicate are not accessed, their values in the resulting vector are undefined.

Definition at line 8433 of file memory.hpp.

gather() [19/28]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> && detail::is_simd_view_type_v<OffsetSimdViewT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		OffsetSimdViewT	byte_offsets,
		PropertyListT	props = {}
	)

simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-ga-9) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 602 of file memory.hpp.

gather() [20/28]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> && detail::is_simd_view_type_v<OffsetSimdViewT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-8) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element of the returned vector is undefined.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read. Elements in masked out lanes are undefined.

Definition at line 573 of file memory.hpp.

gather() [21/28]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> && detail::is_simd_view_type_v<OffsetSimdViewT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-7) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
pass_thru	The vector pass through values.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 538 of file memory.hpp.

gather() [22/28]

template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-ga-6) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 498 of file memory.hpp.

gather() [23/28]

template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N >	byte_offsets,
		simd_mask< N >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (usm-ga-5) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element of the returned vector is undefined.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read. Elements in masked out lanes are undefined.

Definition at line 472 of file memory.hpp.

gather() [24/28]

template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N >	byte_offsets,
		simd_mask< N >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-4) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
pass_thru	The vector pass through values.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 440 of file memory.hpp.

gather() [25/28]

template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 405 of file memory.hpp.

gather() [26/28]

template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements.

Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element of the returned vector is undefined.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read. Elements in masked out lanes are undefined.

Definition at line 348 of file memory.hpp.

gather() [27/28]

template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	const T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-1) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3)

The next 3 functions are similar to the above and were added for convenience. They assume the VS parameter is set to 1 and do not require specifying the template parameters <T, N, VS> at function calls. template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-4) simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (usm-ga-5) simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-ga-6)

The next 3 functions are variations of the first 3 above (usm-ga-1,2,3) and were added only to support simd_view instead of simd for byte_offsets and/or pass_thru operands. template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-7) simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-8) simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-ga-9) template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-1) Supported platforms: DG2, PVC only - Temporary restriction for the variant with pass_thru operand. Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
pass_thru	The vector pass through values.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Returns

A vector of elements read.

Definition at line 289 of file memory.hpp.

gather() [28/28]

template<typename Tx , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, simd<Tx, N> > sycl::_V1::ext::intel::esimd::gather	(	const Tx *	p,
		Toffset	offset,
		simd_mask< N >	mask = 1
	)

A variation of gather API with offsets represented as scalar.

Template Parameters

Tx	Element type, must be of size 4 or less.
N	Number of elements to read; can be 1, 2, 4, 8, 16 or 32.

Parameters

p	The base address.
offset	the scalar 32-bit or 64-bit offset in bytes. ((byte*)p + offset) must be element size aligned.
mask	The access mask, defaults to all 1s.

Returns

A vector of elements read. Elements in masked out lanes are undefined.

Definition at line 620 of file memory.hpp.

gather_rgba() [1/4]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>

__ESIMD_API std::enable_if_t<((N == 8 || N == 16 || N == 32) && sizeof(T) == 4 && detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read>), simd<T, N * get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba	(	AccessorT	acc,
		simd< detail::DeviceAccessorOffsetT, N >	offsets,
		detail::DeviceAccessorOffsetT	global_offset = 0,
		simd_mask< N >	mask = 1
	)

Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets.

Up to 4 32-bit data elements may be accessed at each address depending on the channel mask RGBAMask. Each pixel's address must be 4-byte aligned. For usage examples, see usm_gather_rgba above, the only difference would be the usage of an accessor instead of a usm pointer.

Template Parameters

RGBAMask	A pixel's channel mask.
AccessorT	The accessor type for the memory to be loaded/gathered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.

Parameters

acc	The accessor representing memory address of the access.
offsets	Byte offsets of the pixels relative to the base pointer.
global_offset	Byte offset of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Returns

Read data - up to N*4 values of type Tx.

Definition at line 4100 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), and sycl::_V1::ext::intel::esimd::get_surface_index().

gather_rgba() [2/4]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >

__ESIMD_API std::enable_if_t<detail::is_simd_view_type_v<OffsetSimdViewT>, simd<T, N * get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba	(	const T *	p,
		OffsetSimdViewT	offsets,
		simd_mask< N >	mask = 1
	)

A variation of gather_rgba API with offsets represented as simd_view object.

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
Mask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
offsets	simd_view of byte offsets of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Returns

Read data - up to N*4 values of type Tx.

Definition at line 3938 of file memory.hpp.

gather_rgba() [3/4]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >

__ESIMD_API simd<T, N * get_num_channels_enabled(RGBAMask)> sycl::_V1::ext::intel::esimd::gather_rgba	(	const T *	p,
		simd< Toffset, N >	offsets,
		simd_mask< N >	mask = 1
	)

Gather and transpose pixels from given memory locations defined by the base pointer p and offsets.

Up to 4 32-bit data elements may be accessed at each address depending on the channel mask Mask template parameter. Each pixel's address must be 4 byte aligned. As an example, let's assume we want to read n pixels at address addr, skipping G and B channels. Each channel is a 32-bit float and the pixel data at given address in memory is:

R1 G1 B1 A1 R2 G2 B2 A2 ... Rn Gn Bn An

Then this can be achieved by using

simd<uint32_t, n> byte_offsets(0, 4*4 /* byte size of a single pixel */);
auto x = gather_rgba<float, n, rgba_channel_mask::AR>(addr, byte_offsets);

Returned x will contain 2*n float elements:

R1 R2 ... Rn A1 A2 ... An

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
Mask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
offsets	vector of byte offsets of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Returns

Read data - up to N*4 values of type Tx.

Definition at line 3908 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), and simd_mask< _Tp, _Abi >::data().

gather_rgba() [4/4]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, simd<T, N * get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba	(	const T *	p,
		Toffset	offset,
		simd_mask< N >	mask = 1
	)

A variation of gather_rgba API with offsets represented as scalar.

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
Mask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
offset	scalar byte offsets of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Returns

Read data - up to N*4 values of type Tx.

Definition at line 3961 of file memory.hpp.

get_surface_index()

template<typename AccessorTy >

__ESIMD_API SurfaceIndex sycl::_V1::ext::intel::esimd::get_surface_index

(

AccessorTy

acc

)

Get surface index corresponding to a SYCL accessor.

Parameters

acc	a SYCL buffer or image accessor.

Returns

the index of the corresponding surface (aka "binding table index").

Definition at line 53 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::SLM_BTI.

Referenced by sycl::_V1::ext::intel::esimd::detail::atomic_update_impl(), sycl::_V1::ext::intel::esimd::detail::block_load_impl(), sycl::_V1::ext::intel::esimd::detail::block_store_impl(), sycl::_V1::ext::intel::esimd::gather_rgba(), sycl::_V1::ext::intel::esimd::media_block_load(), sycl::_V1::ext::intel::esimd::media_block_store(), sycl::_V1::ext::intel::esimd::scatter_rgba(), sycl::_V1::ext::intel::esimd::slm_atomic_update(), sycl::_V1::ext::intel::esimd::slm_gather_rgba(), and sycl::_V1::ext::intel::esimd::slm_scatter_rgba().

load_2d()

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = oneapi::experimental::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N> > sycl::_V1::ext::intel::esimd::load_2d	(	const T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y,
		PropertyListT	props = {}
	)

template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = empty_properties_t> simd<T, N> load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block load.

Supported platforms: PVC VISA instruction: lsc_load_block2d.ugm

Collects elements located at specified address and returns them as a single simd object.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
NBlocks	is the number of blocks.
Transposed	is the transposed version or not.
Transformed	is apply VNNI transform or not.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.
props	The optional compile-time properties. Only cache hint properties are used.

Returns

is a vector of type T and size N, where N is BlockWidth * BlockHeight * NBlocks, if not transformed; otherwise, N = roundUpNextMultiple(BlockHeight, 4 / sizeof(T)) * getNextPowerOf2(BlockWidth) * NBlocks

Definition at line 9375 of file memory.hpp.

media_block_load()

template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>

__ESIMD_API simd<T, m * N> sycl::_V1::ext::intel::esimd::media_block_load	(	AccessorTy	acc,
		unsigned	x,
		unsigned	y
	)

Media block load.

Template Parameters

T	is the element data type.
m	is the height of the 2D block.
N	is the width of the 2D block.
AccessorTy	is type of the SYCL accessor.
plane	is planar surface index.

Parameters

acc	is the SYCL accessor.
x	is X-coordinate of the left upper rectangle corner in BYTES.
y	is Y-coordinate of the left upper rectangle corner in ROWS.

Returns

the linearized 2D block data read from surface.

Definition at line 7973 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::get_surface_index().

media_block_store()

template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>

__ESIMD_API void sycl::_V1::ext::intel::esimd::media_block_store	(	AccessorTy	acc,
		unsigned	x,
		unsigned	y,
		simd< T, m *N >	vals
	)

Media block store.

Template Parameters

T	is the element data type.
m	is the height of the 2D block.
N	is the width of the 2D block.
is	AccessorTy type of the SYCL accessor.
plane	is planar surface index.

Parameters

acc	is the SYCL accessor.
x	is X-coordinate of the left upper rectangle corner in BYTES.
y	is Y-coordinate of the left upper rectangle corner in ROWS.
vals	is the linearized 2D block data to be written to surface.

Definition at line 8014 of file memory.hpp.

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

prefetch() [1/20]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (acc-pf-6) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets to the cache.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9209 of file memory.hpp.

prefetch() [2/20]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-5) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (acc + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9177 of file memory.hpp.

prefetch() [3/20]

template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< std::is_integral_v<OffsetT> && detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		OffsetT	byte_offset,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, PropertyListT props = {}); // (acc-pf-8) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the accessor acc, and offset byte_offset and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

acc	Accessor referencing the data to load.
byte_offset	offset from the base address
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9272 of file memory.hpp.

prefetch() [4/20]

template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< std::is_integral_v<OffsetT> && detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		OffsetT	byte_offset,
		simd_mask< 1 >	mask,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (acc-pf-7) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the accessor acc, and offset byte_offset and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

acc	Accessor referencing the data to load.
byte_offset	offset from the base address.
mask	The access mask. If it is set to 0, then the prefetch is omitted.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9239 of file memory.hpp.

prefetch() [5/20]

template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, PropertyListT props = {}); // (acc-pf-10) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the accessor acc and the length VS into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

acc	Accessor referencing the data to load.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9323 of file memory.hpp.

Referenced by sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d().

prefetch() [6/20]

template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		simd< OffsetT, N >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (acc-pf-4) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, into the cache.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9141 of file memory.hpp.

prefetch() [7/20]

template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		simd< OffsetT, N >	byte_offsets,
		simd_mask< N >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (acc-pf-3) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the prefetch from (acc + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9113 of file memory.hpp.

prefetch() [8/20]

template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, into the cache.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9080 of file memory.hpp.

prefetch() [9/20]

template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-1) void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2)

The next 2 functions are similar to the above and were added for convenience. They assume the VS parameter is set to 1 and do not require specifying the template parameters <T, N, VS> at function calls. template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (acc-pf-3) void prefetch(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-pf-4) The next 2 functions are variations of the first 2 above (acc-pf-1,2) and were added only to support simd_view instead of simd for byte_offsets operand. template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {});//(acc-pf-5) void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); //(acc-pf-6)

The next functions perform transposed 1-channel prefetch. template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (acc-pf-7) void prefetch(AccessorT acc, OffsetT byte_offset, PropertyListT props = {}); // (acc-pf-8) template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd_mask<1> mask, PropertyListT props = {}); // (acc-pf-9) void prefetch(AccessorT acc, PropertyListT props = {}); // (acc-pf-10) template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-1) Supported platforms: DG2, PVC only. Prefetches elements of the type 'T' from memory locations addressed by the accessor acc and byte offsets byte_offsets, to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the prefetch from (acc + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

acc	Accessor referencing the data to load.
byte_offsets	the vector of offsets in bytes. If force stateless memory is used the offsets can be up to 64 bit size, otherwise up to 32 bit size. For each i, (acc + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9044 of file memory.hpp.

prefetch() [10/20]

template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	AccessorT	acc,
		simd_mask< 1 >	mask,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd_mask<1> mask, PropertyListT props = {}); //(acc-pf-9) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the accessor acc and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

acc	Accessor referencing the data to load.
mask	The access mask. If it is set to 0, then the prefetch is omitted.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9300 of file memory.hpp.

prefetch() [11/20]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		OffsetSimdViewT	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-pf-6) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets to the cache.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8867 of file memory.hpp.

prefetch() [12/20]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-5) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (p + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8839 of file memory.hpp.

prefetch() [13/20]

template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< std::is_integral_v<OffsetT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		OffsetT	byte_offset,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}); // (usm-pf-8) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the base pointer p, and offset byte_offset and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

p	The base address.
byte_offset	offset from the base address
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8921 of file memory.hpp.

prefetch() [14/20]

template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< std::is_integral_v<OffsetT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		OffsetT	byte_offset,
		simd_mask< 1 >	mask,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-7) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the base pointer p, and offset byte_offset and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

p	The base address.
byte_offset	offset from the base address.
mask	The access mask. If it is set to 0, then the prefetch is omitted.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8894 of file memory.hpp.

prefetch() [15/20]

template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, PropertyListT props = {}); // (usm-pf-10) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the base pointer p and the length VS into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

p	The base address.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8968 of file memory.hpp.

prefetch() [16/20]

template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		simd< OffsetT, N >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-pf-4) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, into the cache.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8808 of file memory.hpp.

prefetch() [17/20]

template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		simd< OffsetT, N >	byte_offsets,
		simd_mask< N >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (usm-pf-3) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the prefetch from (p + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8783 of file memory.hpp.

prefetch() [18/20]

template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, into the cache.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8753 of file memory.hpp.

prefetch() [19/20]

template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-1) void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2)

The next 2 functions are similar to the above and were added for convenience. They assume the VS parameter is set to 1 and do not require specifying the template parameters <T, N, VS> at function calls. template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (usm-pf-3) void prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-pf-4) The next 2 functions are variations of the first 2 above (usm-pf-1,2) and were added only to support simd_view instead of simd for byte_offsets operand. template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-5) void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-pf-6)

The next functions perform transposed 1-channel prefetch. template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-7) void prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}); // (usm-pf-8) template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-9) void prefetch(const T *p, PropertyListT props = {}); //(usm-pf-10) template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-1) Supported platforms: DG2, PVC only. Prefetches elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, to the cache. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the prefetch from (p + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to read.
VS	Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned.
mask	The access mask.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8725 of file memory.hpp.

prefetch() [20/20]

template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch	(	const T *	p,
		simd_mask< 1 >	mask,
		PropertyListT	props = {}
	)

template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}); //(usm-pf-9) Supported platforms: DG2, PVC only.

Prefetches elements of the type 'T' from continuous memory location addressed by the base pointer p and the length VS elements into the cache.

Template Parameters

T	Element type.
VS	Vector size. It specifies the number of consequent elements to prefetch.

Parameters

p	The base address.
mask	The access mask. If it is set to 0, then the prefetch is omitted.
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 8947 of file memory.hpp.

prefetch_2d()

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::prefetch_2d	(	const T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y,
		PropertyListT	props = {}
	)

template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>(), typename PropertyListT = empty_properties_t> void prefetch_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block prefetch.

Supported platforms: PVC VISA instruction: lsc_load_block2d.ugm

Prefetches elements located at specified address.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
NBlocks	is the number of blocks.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.
props	The compile-time properties. Only cache hint properties are used.

Definition at line 9418 of file memory.hpp.

scalar_load()

template<typename T , typename AccessorTy >

__ESIMD_API T sycl::_V1::ext::intel::esimd::scalar_load	(	AccessorTy	acc,
		detail::DeviceAccessorOffsetT	offset
	)

Load a scalar value from an accessor.

Template Parameters

T	Type of the value.
AccessorTy	Type of the accessor.

Parameters

acc	Accessor to load from.
offset	Offset in bytes.

Returns

The loaded value.

Definition at line 3851 of file memory.hpp.

scalar_store()

template<typename T , typename AccessorTy >

__ESIMD_API void sycl::_V1::ext::intel::esimd::scalar_store	(	AccessorTy	acc,
		detail::DeviceAccessorOffsetT	offset,
		T	val
	)

Store a scalar value into an accessor.

Template Parameters

T	Type of the value.
AccessorTy	Type of the accessor.

Parameters

acc	Accessor to store to.
offset	Offset in bytes.
val	The stored value.

Definition at line 3866 of file memory.hpp.

scatter() [1/15]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-4)

Writes ("scatters") elements of the input vector to memory locations addressed by the local accessor acc and byte offsets byte_offsets.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	The accessor to scatter to.
byte_offsets	the vector of 32-bit offsets in bytes represented as a 'simd_view' object. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
props	The optional compile-time properties. Only 'alignment' property is used.

Definition at line 8623 of file memory.hpp.

scatter() [2/15]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-3)

Writes ("scatters") elements of the input vector to memory locations addressed by the local accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input mask.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	The accessor to scatter to.
byte_offsets	the vector of 32-bit offsets in bytes represented as a 'simd_view' object. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 8585 of file memory.hpp.

scatter() [3/15]

template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_write> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorT	acc,
		simd< uint32_t, N/VS >	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-2)

Writes ("scatters") elements of the input vector to memory locations addressed by the local accessor acc and byte offsets byte_offsets.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	The accessor to scatter to.
byte_offsets	the vector of 32-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
props	The optional compile-time properties. Only 'alignment' property is used.

Definition at line 8543 of file memory.hpp.

scatter() [4/15]

template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_write> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorT	acc,
		simd< uint32_t, N/VS >	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

Variant of scatter that uses local accessor as a parameter template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-1)

template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-2) The next two functions are similar to lacc-sc-{1,2} with the 'byte_offsets' parameter represerented as 'simd_view'. template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-3) template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-4) template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-1)

Writes ("scatters") elements of the input vector to memory locations addressed by the local accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input mask.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	The accessor to scatter to.
byte_offsets	the vector of 32-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' property is used.

Definition at line 8508 of file memory.hpp.

scatter() [5/15]

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) && detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		detail::DeviceAccessorOffsetT	glob_offset,
		simd< T, N >	vals,
		simd_mask< N >	mask = 1
	)

Definition at line 3824 of file memory.hpp.

scatter() [6/15]

template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-4)

Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc and byte offsets byte_offsets.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to store.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 3785 of file memory.hpp.

scatter() [7/15]

template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-3)

Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the store to (acc + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to store.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 3751 of file memory.hpp.

scatter() [8/15]

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) && detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		simd< detail::DeviceAccessorOffsetT, N >	offsets,
		simd< T, N >	vals,
		detail::DeviceAccessorOffsetT	glob_offset,
		simd_mask< N >	mask = 1
	)

Writes elements of a simd object into an accessor at given offsets.

An element can be a 1, 2 or 4-byte value.

Template Parameters

T	Element type; can only be a 1,2,4-byte integer, sycl::half or float.
N	The number of vector elements. Can be 1, 8, 16 or 32.
AccessorTy	The accessor type.

Parameters

acc	The accessor to scatter to.
offsets	Per-element offsets in bytes.
vals	Values to write.
glob_offset	Offset in bytes added to each individual element's offset to compute actual memory access offset for that element.
mask	Memory access mask. Elements with zero corresponding mask's predicate are not accessed.

Definition at line 3812 of file memory.hpp.

scatter() [9/15]

template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-2)

Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc and byte offsets byte_offsets.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to store.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 3711 of file memory.hpp.

scatter() [10/15]

template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		simd< OffsetT, N/VS >	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

Accessor-based scatter.

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-1)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-2) The following two functions are similar to acc-sc-{1,2} with the 'byte_offsets' parameter represented as 'simd_view'. template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {});// (acc-sc-3)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-4)

template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-1)

Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the store to (acc + byte_offsets[i]) is skipped.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

acc	Accessor referencing the data to store.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 3664 of file memory.hpp.

scatter() [11/15]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	T *	p,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (usm-sc-4)

Writes ("scatters") elements of the input vector to different memory locations. Each memory location is base address plus an offset - a value of the corresponding element in the input offset vector.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 819 of file memory.hpp.

scatter() [12/15]

template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	T *	p,
		OffsetSimdViewT	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

Writes ("scatters") elements of the input vector to different memory locations.

Each memory location is base address plus an offset - a value of the corresponding element in the input offset vector. Access to any element's memory location can be disabled via the input mask.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 786 of file memory.hpp.

scatter() [13/15]

template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		simd< T, N >	vals,
		PropertyListT	props = {}
	)

Writes ("scatters") elements of the input vector to different memory locations.

Each memory location is base address plus an offset - a value of the corresponding element in the input offset vector.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 750 of file memory.hpp.

scatter() [14/15]

template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter	(	T *	p,
		simd< OffsetT, N/VS >	byte_offsets,
		simd< T, N >	vals,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-1)

template <typename T, int N, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (usm-sc-2) The next two functions are similar to usm-sc-{1,2} with the 'byte_offsets' parameter represented as 'simd_view'. template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-3) template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (usm-sc-4) template <typename T, int N, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-1)

Writes ("scatters") elements of the input vector to different memory locations. Each memory location is base address plus an offset - a value of the corresponding element in the input offset vector. Access to any element's memory location can be disabled via the input mask.

Template Parameters

T	Element type.
N	Number of elements to write.
VS	Vector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.

Parameters

p	The base address.
byte_offsets	the vector of 32-bit or 64-bit offsets in bytes. For each i, ((byte*)p + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
vals	The vector to scatter.
mask	The access mask.
props	The optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 675 of file memory.hpp.

scatter() [15/15]

template<typename Tx , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && N == 1> sycl::_V1::ext::intel::esimd::scatter	(	Tx *	p,
		Toffset	offset,
		simd< Tx, N >	vals,
		simd_mask< N >	mask = 1
	)

A variation of scatter API with offsets represented as scalar.

Template Parameters

Tx	Element type, must be of size 4 or less.
N	Number of elements to write; can be 1, 2, 4, 8, 16 or 32.

Parameters

p	The base address.
offset	the scalar 32-bit or 64-bit offset in bytes. ((byte*)p + offset) must be element size aligned.
vals	The vector to scatter.
mask	The access mask, defaults to all 1s.

Definition at line 838 of file memory.hpp.

scatter_rgba() [1/5]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>

__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4 && detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::esimd::scatter_rgba	(	AccessorT	acc,
		simd< detail::DeviceAccessorOffsetT, N >	offsets,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		detail::DeviceAccessorOffsetT	global_offset = 0,
		simd_mask< N >	mask = 1
	)

Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector.

See usm_gather_rgba for information about the operation semantics and parameter restrictions/interdependencies.

Template Parameters

RGBAMask	Pixel's channel mask.
AccessorT	The accessor type for the memory to be stored/scattered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
N	The number of elements to access.

Parameters

offsets	Byte offsets of each element.
vals	values to be written.
global_offset	Byte offset of the pixels relative to the base pointer.
mask	Operation mask. All-1 by default.

Definition at line 4154 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), and sycl::_V1::ext::intel::esimd::get_surface_index().

scatter_rgba() [2/5]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >

__ESIMD_API std::enable_if_t<detail::is_simd_view_type_v<OffsetSimdViewT> > sycl::_V1::ext::intel::esimd::scatter_rgba	(	T *	p,
		OffsetSimdViewT	offsets,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		simd_mask< N >	mask = 1
	)

A variation of scatter_rgba API with offsets represented as simd_view object.

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
RGBAMask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
vals	values to be written.
offsets	simd_view of byte offsets of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Definition at line 4030 of file memory.hpp.

scatter_rgba() [3/5]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >

__ESIMD_API void sycl::_V1::ext::intel::esimd::scatter_rgba	(	T *	p,
		simd< Toffset, N >	offsets,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		simd_mask< N >	mask = 1
	)

Transpose and scatter pixels to given memory locations defined by the base pointer p and offsets.

Up to 4 32-bit data elements may be accessed at each address depending on the channel mask RGBAMask. Each pixel's address must be 4 byte aligned. This is basically an inverse operation for gather_rgba. Unlike gather_rgba, this function imposes restrictions on possible Mask template argument values. It can only be one of the following: ABGR, BGR, GR, R.

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
RGBAMask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
vals	values to be written.
offsets	vector of byte offsets of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Definition at line 3998 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), and simd_mask< _Tp, _Abi >::data().

scatter_rgba() [4/5]

template<typename T , int N, rgba_channel_mask RGBAMask>

__ESIMD_API std:: enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4> sycl::_V1::ext::intel::esimd::scatter_rgba	(	T *	p,
		simd< uint32_t, N >	offsets,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		simd_mask< N >	mask = 1
	)

Definition at line 4063 of file memory.hpp.

scatter_rgba() [5/5]

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && N == 1> sycl::_V1::ext::intel::esimd::scatter_rgba	(	T *	p,
		Toffset	offset,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		simd_mask< N >	mask = 1
	)

A variation of scatter_rgba API with offsets represented as scalar.

Template Parameters

T	Element type of the returned vector. Must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
RGBAMask	A pixel's channel mask.

Parameters

p	The USM base pointer representing memory address of the access.
vals	values to be written.
offset	scalar byte offset of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Definition at line 4054 of file memory.hpp.

split_barrier() [1/2]

template<split_barrier_action flag>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::split_barrier

(

)

Generic work-group split barrier.

Template Parameters

flag	- split barrier action.

Definition at line 28 of file memory.hpp.

split_barrier() [2/2]

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::split_barrier

(

split_barrier_action

flag

)

Definition at line 33 of file memory.hpp.

store_2d()

template<typename T , int BlockWidth, int BlockHeight = 1, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>

__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::store_2d	(	T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y,
		simd< T, N >	Vals,
		PropertyListT	props = {}
	)

2D USM pointer block store.

Supported platforms: PVC VISA instruction: lsc_store_block2d.ugm

Stores elements at specified address.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.
Vals	is a vector to store of type T and size N, where N = BlockWidth * BlockHeight
props	The optional compile-time properties. Only cache hint properties are used.

Definition at line 9454 of file memory.hpp.