ESIMD APIs to access memory via accessors, USM pointers, perform per-element atomic operations. More...
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. | |
Named barrier APIs. | |
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. | |
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::l2_flush_texture_data = 0x4 , sycl::_V1::ext::intel::esimd::l2_flush_constant_data = 0x8 , sycl::_V1::ext::intel::esimd::l2_flush_rw_data = 0x10 , sycl::_V1::ext::intel::esimd::local_barrier = 0x20 , sycl::_V1::ext::intel::esimd::l1_flush_ro_data = 0x40 } |
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::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::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::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::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::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::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::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<int VS = 1, typename OffsetT , typename T , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< PassThruSimdViewT >, simd< T, N > > | sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PassThruSimdViewT pass_thru, PropertyListT props={}) |
template <int VS = 1, typename OffsetT, typename T, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); Variation of the API that allows using simd_view without specifying T and N template parameters. More... | |
template<int VS = 1, typename OffsetSimdViewT , typename T , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&detail::is_simd_view_type_v< PassThruSimdViewT >, simd< T, N > > | sycl::_V1::ext::intel::esimd::gather (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PassThruSimdViewT pass_thru, PropertyListT props={}) |
template <int VS = 1, typename OffsetSimdViewT, typename T, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); Variation of the API that allows using simd_view without specifying T and N template parameters. More... | |
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::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::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::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<int VS = 1, typename OffsetT , typename ValuesSimdViewT , typename T , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, ValuesSimdViewT vals, simd_mask< N/VS > mask, PropertyListT props={}) |
template <int VS = 1, typename OffsetT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {}); More... | |
template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::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={}) |
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) More... | |
template<int VS = 1, typename OffsetSimdViewT , typename ValuesSimdViewT , typename T , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< ValuesSimdViewT > &&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, ValuesSimdViewT vals, simd_mask< N/VS > mask, PropertyListT props={}) |
template <int VS = 1, typename OffsetSimdViewT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {}); More... | |
template<int VS = 1, typename OffsetT , typename ValuesSimdViewT , typename T , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, ValuesSimdViewT vals, PropertyListT props={}) |
template <int VS = 1, typename OffsetT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, PropertyListT props = {}); More... | |
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::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={}) |
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) More... | |
template<int VS, typename OffsetSimdViewT , typename T , int N, typename PropertyListT = ext::oneapi::experimental::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 <int VS, typename OffsetSimdViewT, typename T, int N, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T,N> vals, PropertyListT props = {}); More... | |
template<int VS = 1, typename OffsetSimdViewT , typename ValuesSimdViewT , typename T , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&detail::is_simd_view_type_v< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (T *p, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props={}) |
template <int VS = 1, typename OffsetSimdViewT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props = {}); 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::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::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::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::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::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::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::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<int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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 > &&detail::is_simd_view_type_v< PassThruSimdViewT > &&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, PassThruSimdViewT pass_thru, PropertyListT props={}) |
template <int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>> simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets and pass_thru are represented as simd_view . More... | |
template<int VS = 1, typename AccessorT , typename OffsetT , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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< PassThruSimdViewT > &&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, PassThruSimdViewT pass_thru, PropertyListT props={}) |
template <int VS = 1, typename AccessorT, typename OffsetT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>> simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets is represented as simd_view . More... | |
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::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::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::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::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::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<int VS, typename AccessorTy , typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::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 <int VS, typename AccessorTy, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); More... | |
template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetSimdViewT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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 > &&detail::is_simd_view_type_v< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask< N/VS > mask, PropertyListT props={}) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {}); More... | |
template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetSimdViewT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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 > &&detail::is_simd_view_type_v< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props={}) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props = {}); More... | |
template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, ValuesSimdViewT vals, simd_mask< N/VS > mask, PropertyListT props={}) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {}); More... | |
template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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< ValuesSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > > | sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, ValuesSimdViewT vals, PropertyListT props={}) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, PropertyListT props = {}); 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<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... | |
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::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::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::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::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::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::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<int VS = 1, typename AccessorT , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t> | |
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< PassThruSimdViewT > &&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, PassThruSimdViewT pass_thru, PropertyListT props={}) |
template <int VS = 1, typename AccessorT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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 = {}); This function is identical to (lacc-ga-1) except that the pass_thru is represented as simd_view . 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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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::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 APIs to access memory via accessors, USM pointers, perform per-element atomic operations.
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.
Definition at line 12056 of file memory.hpp.
__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 12106 of file memory.hpp.
References sycl::_V1::ext::intel::esimd::global_coherent_fence, and sycl::_V1::ext::intel::esimd::local_barrier.
__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.
T | Element type. |
N | Number of elements to load. |
AccessorTy | Accessor type (auto-deduced). |
Flags | The alignment specifier type tag. |
acc | The local accessor. |
byte_offset | The offset to load from in bytes. |
Flags | Specifies the alignment. |
Definition at line 12224 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to store. |
AccessorT | Accessor type (auto-deduced). |
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 12251 of file memory.hpp.
__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.
Tx | Element type. |
N | Number of elements to store. |
Flags | The alignment specifier type tag. |
addr | The memory address to store at. |
vals | The vector to store. |
Flags | Specifies the alignment. |
Definition at line 1703 of file memory.hpp.
References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data().
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence | ( | ) |
esimd::fence sets the memory read/write order.
cntl | A bitmask composed from fence_mask bits. |
Definition at line 12076 of file memory.hpp.
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence | ( | ) |
Memory fence.
Supported platforms: DG2, PVC
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 12087 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.
__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.
T | Element type. |
N | The number of vector elements. |
AccessorT | The accessor type. |
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 4326 of file memory.hpp.
__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 4750 of file memory.hpp.
__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> && detail::is_simd_view_type_v<PassThruSimdViewT> && 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, | ||
PassThruSimdViewT | pass_thru, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>> simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets
and pass_thru
are represented as simd_view
.
Definition at line 4674 of file memory.hpp.
__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 4729 of file memory.hpp.
__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,
template <int VS, typename T, int N, 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
.
simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets
is represented as simd_view
.
Definition at line 4621 of file memory.hpp.
__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 |
||
) |
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.
T | Element type. |
N | The number of vector elements. |
AccessorT | The accessor type. |
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 4288 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().
__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 4580 of file memory.hpp.
__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 4555 of file memory.hpp.
__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 4600 of file memory.hpp.
__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.
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. |
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. |
Definition at line 4527 of file memory.hpp.
__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<PassThruSimdViewT> && 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, | ||
PassThruSimdViewT | pass_thru, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorT, typename OffsetT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>> simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets
is represented as simd_view
.
Definition at line 4706 of file memory.hpp.
__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.
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. |
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. |
Definition at line 4473 of file memory.hpp.
__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.
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. |
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. |
Definition at line 4424 of file memory.hpp.
__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 12473 of file memory.hpp.
__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 12446 of file memory.hpp.
__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 12494 of file memory.hpp.
__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.
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. |
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. |
Definition at line 12418 of file memory.hpp.
__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<PassThruSimdViewT> && 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, | ||
PassThruSimdViewT | pass_thru, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::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 = {}); This function is identical to (lacc-ga-1) except that the pass_thru
is represented as simd_view
.
Definition at line 12601 of file memory.hpp.
__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.
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. |
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. |
Definition at line 12380 of file memory.hpp.
__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.
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. |
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. |
Definition at line 12335 of file memory.hpp.
__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.
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. |
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 12671 of file memory.hpp.
__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.
simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); Variation of the API that allows using simd_view
without specifying T
and N
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. |
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. |
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.
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. |
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. |
Definition at line 769 of file memory.hpp.
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> && detail::is_simd_view_type_v<OffsetSimdViewT> && detail::is_simd_view_type_v<PassThruSimdViewT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather | ( | const T * | p, |
OffsetSimdViewT | byte_offsets, | ||
simd_mask< N/VS > | mask, | ||
PassThruSimdViewT | pass_thru, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetSimdViewT, typename T, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); Variation of the API that allows using simd_view
without specifying T
and N
template parameters.
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.
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. |
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. |
Definition at line 624 of file memory.hpp.
__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.
simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); Variation of the API that allows using simd_view
without specifying T
and N
template parameters.
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.
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. |
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. |
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.
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. |
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. |
Definition at line 706 of file memory.hpp.
__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.
template <int VS = 1, typename OffsetSimdViewT, typename T, int N, 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 = {}); Variation of the API that allows using simd_view
without specifying T
and N
template parameters.
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.
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. |
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. |
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.
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. |
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. |
Definition at line 540 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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. |
Definition at line 500 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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. |
Definition at line 474 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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. |
Definition at line 442 of file memory.hpp.
__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.
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. |
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. |
Definition at line 407 of file memory.hpp.
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v<PropertyListT> && detail::is_simd_view_type_v<PassThruSimdViewT>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather | ( | const T * | p, |
simd< OffsetT, N/VS > | byte_offsets, | ||
simd_mask< N/VS > | mask, | ||
PassThruSimdViewT | pass_thru, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetT, typename T, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); Variation of the API that allows using simd_view
without specifying T
and N
template parameters.
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.
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. |
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. |
Definition at line 581 of file memory.hpp.
__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.
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. |
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. |
Definition at line 348 of file memory.hpp.
__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.
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. |
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. |
Definition at line 289 of file memory.hpp.
__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.
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 . |
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. |
Definition at line 814 of file memory.hpp.
__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.
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. |
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. |
Tx
. Definition at line 5492 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().
__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.
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. |
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. |
Tx
. Definition at line 5340 of file memory.hpp.
__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:
Then this can be achieved by using
Returned x
will contain 2*n
float
elements:
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. |
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. |
Tx
. Definition at line 5310 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().
__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.
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. |
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. |
Tx
. Definition at line 5363 of file memory.hpp.
__ESIMD_API SurfaceIndex sycl::_V1::ext::intel::esimd::get_surface_index | ( | AccessorTy | acc | ) |
Get surface index corresponding to a SYCL accessor.
acc | a SYCL buffer or image accessor. |
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().
__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.
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 |
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. |
Definition at line 13940 of file memory.hpp.
__ESIMD_API simd<T, m * N> sycl::_V1::ext::intel::esimd::media_block_load | ( | AccessorTy | acc, |
unsigned | x, | ||
unsigned | y | ||
) |
Media block load.
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. |
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. |
Definition at line 12128 of file memory.hpp.
References sycl::_V1::ext::intel::esimd::get_surface_index(), and sycl::_V1::ext::intel::esimd::detail::isPowerOf2().
__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.
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. |
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 12170 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().
__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.
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'. |
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 13766 of file memory.hpp.
__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.
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'. |
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 13735 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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 13837 of file memory.hpp.
__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. The maximum size of prefetched block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed VS
values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed VS
values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed VS
values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed VS
values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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. |
Definition at line 13804 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
acc | Accessor referencing the data to load. |
props | The optional compile-time properties. Only cache hint properties are used. |
Definition at line 13888 of file memory.hpp.
Referenced by sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d().
__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.
T | Element type. |
N | Number of elements to read. |
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 13700 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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 13672 of file memory.hpp.
__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.
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'. |
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 13639 of file memory.hpp.
__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.
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'. |
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 13603 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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 13865 of file memory.hpp.
__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.
template <int VS = 1, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); 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.
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'. |
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. |
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. Prefetches elements of the type 'T' from memory locations addressed by the base pointer p
and byte offsets byte_offsets
to the cache.
VS | Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. |
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 13357 of file memory.hpp.
__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.
template <int VS = 1, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); 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.
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'. |
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. |
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
VS | Vector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. |
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 13329 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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 13480 of file memory.hpp.
__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. The maximum size of a prefetched block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) is equal to 8 the address must be 8-byte aligned. Also, 8-byte alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed VS
values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed VS
values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed VS
values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed VS
values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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. |
Definition at line 13453 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
p | The base address. |
props | The optional compile-time properties. Only cache hint properties are used. |
Definition at line 13527 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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 13298 of file memory.hpp.
__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.
T | Element type. |
N | Number of elements to read. |
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 13273 of file memory.hpp.
__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.
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'. |
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 13243 of file memory.hpp.
__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.
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'. |
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 13215 of file memory.hpp.
__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.
T | Element type. |
VS | Vector size. It specifies the number of consequent elements to prefetch. |
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 13506 of file memory.hpp.
__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.
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 |
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 13983 of file memory.hpp.
__ESIMD_API T sycl::_V1::ext::intel::esimd::scalar_load | ( | AccessorTy | acc, |
detail::DeviceAccessorOffsetT | offset | ||
) |
Load a scalar value from an accessor.
T | Type of the value. |
AccessorTy | Type of the accessor. |
acc | Accessor to load from. |
offset | Offset in bytes. |
Definition at line 5253 of file memory.hpp.
__ESIMD_API void sycl::_V1::ext::intel::esimd::scalar_store | ( | AccessorTy | acc, |
detail::DeviceAccessorOffsetT | offset, | ||
T | val | ||
) |
Store a scalar value into an accessor.
T | Type of the value. |
AccessorTy | Type of the accessor. |
acc | Accessor to store to. |
offset | Offset in bytes. |
val | The stored value. |
Definition at line 5268 of file memory.hpp.
__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
.
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. |
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 13113 of file memory.hpp.
__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.
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. |
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 12822 of file memory.hpp.
__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
.
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. |
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 12781 of file memory.hpp.
__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.
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. |
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 12746 of file memory.hpp.
__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 5226 of file memory.hpp.
__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 <int VS, typename AccessorTy, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void 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)
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc
and byte offsets byte_offsets
.
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. |
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. |
Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc
and byte offsets byte_offsets
.
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. |
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 4974 of file memory.hpp.
__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)
template <int VS, typename AccessorTy, typename T, int N, 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 = {});
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.
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. |
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. |
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 4899 of file memory.hpp.
__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> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | AccessorTy | acc, |
OffsetSimdViewT | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc
and byte offsets byte_offsets
.
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. |
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 5067 of file memory.hpp.
__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> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | AccessorTy | acc, |
OffsetSimdViewT | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
simd_mask< N/VS > | mask, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 5023 of file memory.hpp.
__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.
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. |
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 5214 of file memory.hpp.
__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
.
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. |
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 4859 of file memory.hpp.
__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 = {} |
||
) |
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.
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. |
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 4812 of file memory.hpp.
__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<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | AccessorTy | acc, |
simd< OffsetT, N/VS > | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the accessor acc
and byte offsets byte_offsets
.
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. |
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 5153 of file memory.hpp.
__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<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | AccessorTy | acc, |
simd< OffsetT, N/VS > | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
simd_mask< N/VS > | mask, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 5114 of file memory.hpp.
__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 <int VS, typename OffsetSimdViewT, typename T, int N, typename PropertyListT = empty_properties_t> void 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)
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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. |
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.
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. |
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 1167 of file memory.hpp.
__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 = {} |
||
) |
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 <int VS, typename OffsetSimdViewT, typename T, int N, typename PropertyListT = empty_properties_t> void 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.
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. |
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. |
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 1094 of file memory.hpp.
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<OffsetSimdViewT> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | T * | p, |
OffsetSimdViewT | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetSimdViewT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 1242 of file memory.hpp.
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<ValuesSimdViewT> && 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, | ||
ValuesSimdViewT | vals, | ||
simd_mask< N/VS > | mask, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetSimdViewT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 1021 of file memory.hpp.
__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 = {} |
||
) |
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)
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.
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. |
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 984 of file memory.hpp.
__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.
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. |
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 869 of file memory.hpp.
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | T * | p, |
simd< OffsetT, N/VS > | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 1059 of file memory.hpp.
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter | ( | T * | p, |
simd< OffsetT, N/VS > | byte_offsets, | ||
ValuesSimdViewT | vals, | ||
simd_mask< N/VS > | mask, | ||
PropertyListT | props = {} |
||
) |
template <int VS = 1, typename OffsetT, typename ValuesSimdViewT, typename T, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});
Variation of the API that allows using simd_view
without specifying T
and N
template parameters. 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.
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. |
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 953 of file memory.hpp.
__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.
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 . |
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 1264 of file memory.hpp.
__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.
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. |
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 5546 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().
__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.
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. |
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 5432 of file memory.hpp.
__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
.
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. |
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 5400 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().
__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.
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. |
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 5456 of file memory.hpp.
__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::split_barrier | ( | ) |
Generic work-group split barrier.
flag | - split barrier action. |
Definition at line 29 of file memory.hpp.
__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.
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 |
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 14019 of file memory.hpp.