ESIMD core.

Core APIs defining main vector data types and their interfaces. More...

Collaboration diagram for ESIMD core.:

Modules
	C++ binary operators overloads for ESIMD.
	Standard C++ binary operators overloads applicable to simd_obj_impl derivatives - simd , simd_mask , simd_view and their combinations. The following overloads are defined:
	Alignment control.
	Alignment type tags and related APIs for use with ESIMD memory access operations. The basic restrictions for memory location specified as parameters for memory access APIs supporting alignment control are as follows:
	Main vector data types.
	ESIMD defines the following two main vector data types:

Namespaces
	sycl::_V1::ext::intel::esimd::detail
	sycl::_V1::ext::intel::experimental::esimd::detail

Classes
struct	sycl::_V1::ext::intel::esimd::saturation_on_tag
	Gen hardware supports applying saturation to results of certain operations. More...
struct	sycl::_V1::ext::intel::esimd::saturation_off_tag
	This type tag represents "saturation off" behavior. More...

Macros
#define	__ESIMD_USM_DWORD_ATOMIC_TO_LSC

Typedefs
using	sycl::_V1::ext::intel::esimd::uchar = unsigned char
using	sycl::_V1::ext::intel::esimd::ushort = unsigned short
using	sycl::_V1::ext::intel::esimd::uint = unsigned int
using	sycl::_V1::ext::intel::esimd::SurfaceIndex = unsigned int
	Surface index type. More...

Enumerations
enum	sycl::_V1::ext::intel::esimd::rgba_channel : uint8_t { sycl::_V1::ext::intel::esimd::rgba_channel::R, sycl::_V1::ext::intel::esimd::rgba_channel::G, sycl::_V1::ext::intel::esimd::rgba_channel::B, sycl::_V1::ext::intel::esimd::rgba_channel::A }
	Represents a pixel's channel. More...
enum	sycl::_V1::ext::intel::esimd::rgba_channel_mask : uint8_t {   sycl::_V1::ext::intel::esimd::rgba_channel_mask::R = detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::G = detail::chG, sycl::_V1::ext::intel::esimd::rgba_channel_mask::GR = detail::chG | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::B = detail::chB,   sycl::_V1::ext::intel::esimd::rgba_channel_mask::BR = detail::chB | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::BG = detail::chB | detail::chG, sycl::_V1::ext::intel::esimd::rgba_channel_mask::BGR = detail::chB | detail::chG | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::A = detail::chA,   sycl::_V1::ext::intel::esimd::rgba_channel_mask::AR = detail::chA | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::AG = detail::chA | detail::chG, sycl::_V1::ext::intel::esimd::rgba_channel_mask::AGR = detail::chA | detail::chG | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::AB = detail::chA | detail::chB,   sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABR = detail::chA | detail::chB | detail::chR, sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABG = detail::chA | detail::chB | detail::chG, sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABGR = detail::chA | detail::chB | detail::chG | detail::chR }
	Represents a pixel's channel mask - all possible combinations of enabled channels. More...
enum	sycl::_V1::ext::intel::esimd::atomic_op : uint8_t {   sycl::_V1::ext::intel::esimd::atomic_op::add = 0x0, sycl::_V1::ext::intel::esimd::atomic_op::sub = 0x1, sycl::_V1::ext::intel::esimd::atomic_op::inc = 0x2, sycl::_V1::ext::intel::esimd::atomic_op::dec = 0x3,   sycl::_V1::ext::intel::esimd::atomic_op::umin = 0x4, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("use umin") = umin, sycl::_V1::ext::intel::esimd::atomic_op::umax = 0x5, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("use smax") = umax,   sycl::_V1::ext::intel::esimd::atomic_op::xchg = 0x6, sycl::_V1::ext::intel::esimd::atomic_op::cmpxchg = 0x7, sycl::_V1::ext::intel::esimd::atomic_op::bit_and = 0x8, sycl::_V1::ext::intel::esimd::atomic_op::bit_or = 0x9,   sycl::_V1::ext::intel::esimd::atomic_op::bit_xor = 0xa, sycl::_V1::ext::intel::esimd::atomic_op::smin = 0xb, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("use smin") = smin, sycl::_V1::ext::intel::esimd::atomic_op::smax = 0xc,   sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("use smax") = 0xc, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("fmax" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x10, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("fmin" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x11, sycl::_V1::ext::intel::esimd::atomic_op::fcmpxchg = 0x12,   sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("fcmpwr" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = fcmpxchg, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("fadd" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x13, sycl::_V1::ext::intel::esimd::atomic_op::__SYCL_DEPRECATED =("fsub" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x14, sycl::_V1::ext::intel::esimd::atomic_op::load = 0x15,   sycl::_V1::ext::intel::esimd::atomic_op::store = 0x16, sycl::_V1::ext::intel::esimd::atomic_op::predec = 0xff }
	Represents an atomic operation. More...
enum	sycl::_V1::ext::intel::experimental::esimd::lsc_scope : uint8_t {   sycl::_V1::ext::intel::experimental::esimd::lsc_scope::group = 0, sycl::_V1::ext::intel::experimental::esimd::lsc_scope::local = 1, sycl::_V1::ext::intel::experimental::esimd::lsc_scope::tile = 2, sycl::_V1::ext::intel::experimental::esimd::lsc_scope::gpu = 3,   sycl::_V1::ext::intel::experimental::esimd::lsc_scope::gpus = 4, sycl::_V1::ext::intel::experimental::esimd::lsc_scope::system = 5, sycl::_V1::ext::intel::experimental::esimd::lsc_scope::sysacq = 6 }
	The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC. More...
enum	sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op : uint8_t {   sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::none = 0, sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::evict = 1, sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::invalidate = 2, sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::discard = 3,   sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::clean = 4, sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op::flushl3 = 5 }
	The lsc_fence operation to apply to caches Supported platforms: DG2, PVC. More...
enum	sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind : uint8_t { sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind::untyped_global = 0, sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind::untyped_global_low_pri = 1, sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind::typed_global = 2, sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind::shared_local = 3 }
	The specific LSC shared function to fence with lsc_fence Supported platforms: DG2, PVC. More...
enum	sycl::_V1::ext::intel::experimental::esimd::lsc_data_size : uint8_t {   sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::default_size = 0, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u8 = 1, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u16 = 2, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u32 = 3,   sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u64 = 4, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u8u32 = 5, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u16u32 = 6, sycl::_V1::ext::intel::experimental::esimd::lsc_data_size::u16u32h = 7 }
	Data size or format to read or store. More...
enum	sycl::_V1::ext::intel::experimental::esimd::cache_hint : uint8_t {   sycl::_V1::ext::intel::experimental::esimd::cache_hint::none = 0, sycl::_V1::ext::intel::experimental::esimd::cache_hint::uncached = 1, sycl::_V1::ext::intel::experimental::esimd::cache_hint::cached = 2, sycl::_V1::ext::intel::experimental::esimd::cache_hint::write_back = 3,   sycl::_V1::ext::intel::experimental::esimd::cache_hint::write_through = 4, sycl::_V1::ext::intel::experimental::esimd::cache_hint::streaming = 5, sycl::_V1::ext::intel::experimental::esimd::cache_hint::read_invalidate = 6 }
	L1 or L3 cache hint kinds. More...
enum	sycl::_V1::ext::intel::experimental::esimd::split_barrier_action : uint8_t { sycl::_V1::ext::intel::experimental::esimd::split_barrier_action::wait = 0, sycl::_V1::ext::intel::experimental::esimd::split_barrier_action::signal = 1 }
	Represents a split barrier action. More...

Functions
constexpr int	sycl::_V1::ext::intel::esimd::is_channel_enabled (rgba_channel_mask M, rgba_channel Ch)
constexpr int	sycl::_V1::ext::intel::esimd::get_num_channels_enabled (rgba_channel_mask M)

Variables
static constexpr saturation_off_tag	sycl::_V1::ext::intel::esimd::saturation_off {}
	Type tag object representing "saturation off" behavior. More...
static constexpr saturation_on_tag	sycl::_V1::ext::intel::esimd::saturation_on {}
	Type tag object representing "saturation on" behavior. More...

Detailed Description

Core APIs defining main vector data types and their interfaces.

Macro Definition Documentation

__ESIMD_USM_DWORD_ATOMIC_TO_LSC

#define __ESIMD_USM_DWORD_ATOMIC_TO_LSC

Value:

  " is supported only on ACM, PVC. USM-based atomic will be auto-converted "   \
  "to LSC version."

Definition at line 135 of file common.hpp.

Typedef Documentation

SurfaceIndex

using sycl::_V1::ext::intel::esimd::SurfaceIndex = typedef unsigned int

Surface index type.

Surface is an internal representation of a memory block addressable by GPU in "stateful" memory model, and each surface is identified by its "binding table index" - surface index.

Definition at line 64 of file common.hpp.

uchar

using sycl::_V1::ext::intel::esimd::uchar = typedef unsigned char

Definition at line 41 of file common.hpp.

uint

using sycl::_V1::ext::intel::esimd::uint = typedef unsigned int

Definition at line 43 of file common.hpp.

ushort

using sycl::_V1::ext::intel::esimd::ushort = typedef unsigned short

Definition at line 42 of file common.hpp.

Enumeration Type Documentation

atomic_op

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

strong

Represents an atomic operation.

Operations always return the old value(s) of the target memory location(s) as it was before the operation was applied. Each operation is annotated with a pseudocode illustrating its semantics, addr is a memory address (one of the many, as the atomic operation is vector) the operation is applied at, src0 is its first argumnet, src1 - second.

Enumerator
add	Addition: *addr = *addr + src0.
sub	Subtraction: *addr = *addr - src0.
inc	Increment: *addr = *addr + 1.
dec	Decrement: *addr = *addr - 1.
umin	Minimum: *addr = min(*addr, src0).
__SYCL_DEPRECATED
umax	Maximum: *addr = max(*addr, src0).
__SYCL_DEPRECATED
xchg	Exchange. *addr == src0;
cmpxchg	Compare and exchange. if (*addr == src0) *sddr = src1;
bit_and	Bit and: *addr = *addr & src0.
bit_or	Bit or: *addr = *addr | src0.
bit_xor	Bit xor: *addr = *addr | src0.
smin	Minimum (signed integer): *addr = min(*addr, src0).
__SYCL_DEPRECATED
smax	Maximum (signed integer): *addr = max(*addr, src0).
__SYCL_DEPRECATED
__SYCL_DEPRECATED	Minimum (floating point): *addr = min(*addr, src0).
__SYCL_DEPRECATED	Maximum (floating point): *addr = max(*addr, src0).
fcmpxchg	Compare and exchange (floating point). if (*addr == src0) *addr = src1;
__SYCL_DEPRECATED
__SYCL_DEPRECATED
__SYCL_DEPRECATED
load
store
predec	Decrement: *addr = *addr - 1. The only operation which returns new value of the destination rather than old.

Definition at line 145 of file common.hpp.

cache_hint

enum sycl::_V1::ext::intel::experimental::esimd::cache_hint : uint8_t

strong

L1 or L3 cache hint kinds.

Enumerator
none
uncached
cached
write_back
write_through
streaming
read_invalidate

Definition at line 219 of file common.hpp.

lsc_data_size

enum sycl::_V1::ext::intel::experimental::esimd::lsc_data_size : uint8_t

strong

Data size or format to read or store.

Enumerator
default_size
u8
u16
u32
u64
u8u32
u16u32	load 8b, zero extend to 32b; store the opposite
u16u32h	load 16b, zero extend to 32b; store the opposite

Definition at line 65 of file common.hpp.

lsc_fence_op

enum sycl::_V1::ext::intel::experimental::esimd::lsc_fence_op : uint8_t

strong

The lsc_fence operation to apply to caches Supported platforms: DG2, PVC.

Enumerator
none
evict	no operation
invalidate	dirty lines evicted and invalidated from L1
discard	invalidate all clean lines
clean	direct and clean lines are discarded w/o eviction
flushl3	dirty lines are written to memory, but retained in cache in clean state

Definition at line 45 of file common.hpp.

lsc_memory_kind

enum sycl::_V1::ext::intel::experimental::esimd::lsc_memory_kind : uint8_t

strong

The specific LSC shared function to fence with lsc_fence Supported platforms: DG2, PVC.

Enumerator
untyped_global
untyped_global_low_pri	untyped global memory
typed_global	low-priority untyped global memory
shared_local	typed global memory

Definition at line 57 of file common.hpp.

lsc_scope

enum sycl::_V1::ext::intel::experimental::esimd::lsc_scope : uint8_t

strong

The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC.

Enumerator
group
local	flush out to the threadgroup's scope
tile	flush out to the local scope
gpu	tile, flush out to several DSSs
gpus	entire GPU, flush out to the GPUs LLC
system	all GPUs in the system, flush out to memory shared by all GPUs
sysacq	the entire system memory space

Definition at line 33 of file common.hpp.

rgba_channel

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

strong

Represents a pixel's channel.

Enumerator
R
G
B
A

Definition at line 59 of file common.hpp.

rgba_channel_mask

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

strong

Represents a pixel's channel mask - all possible combinations of enabled channels.

Enumerator
R
G
GR
B
BR
BG
BGR
A
AR
AG
AGR
AB
ABR
ABG
ABGR

Definition at line 105 of file common.hpp.

split_barrier_action

enum sycl::_V1::ext::intel::experimental::esimd::split_barrier_action : uint8_t

strong

Represents a split barrier action.

Enumerator
wait
signal

Definition at line 297 of file common.hpp.

Function Documentation

get_num_channels_enabled()

constexpr int sycl::_V1::ext::intel::esimd::get_num_channels_enabled ( rgba_channel_mask  M )

constexpr

Definition at line 128 of file common.hpp.

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

is_channel_enabled()

constexpr int sycl::_V1::ext::intel::esimd::is_channel_enabled ( rgba_channel_mask  M, rgba_channel  Ch  )

constexpr

Definition at line 123 of file common.hpp.

Referenced by sycl::_V1::ext::intel::esimd::get_num_channels_enabled().

Variable Documentation

saturation_off

constexpr saturation_off_tag sycl::_V1::ext::intel::esimd::saturation_off {}

inlinestaticconstexpr

Type tag object representing "saturation off" behavior.

Definition at line 53 of file common.hpp.

saturation_on

constexpr saturation_on_tag sycl::_V1::ext::intel::esimd::saturation_on {}

inlinestaticconstexpr

Type tag object representing "saturation on" behavior.

Definition at line 56 of file common.hpp.