DPC++ Runtime
Runtime libraries for oneAPI DPC++

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::umax = 0x5, 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::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 140 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

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

umax 

Maximum: *addr = max(*addr, src0).

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

smax 

Maximum (signed integer): *addr = max(*addr, src0).

__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 150 of file common.hpp.

◆ cache_hint

L1 or L3 cache hint kinds.

Enumerator
none 
uncached 
cached 
write_back 
write_through 
streaming 
read_invalidate 

Definition at line 215 of file common.hpp.

◆ lsc_data_size

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 61 of file common.hpp.

◆ lsc_fence_op

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 41 of file common.hpp.

◆ lsc_memory_kind

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 53 of file common.hpp.

◆ lsc_scope

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 29 of file common.hpp.

◆ rgba_channel

Represents a pixel's channel.

Enumerator

Definition at line 59 of file common.hpp.

◆ rgba_channel_mask

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

Enumerator
GR 
BR 
BG 
BGR 
AR 
AG 
AGR 
AB 
ABR 
ABG 
ABGR 

Definition at line 110 of file common.hpp.

◆ split_barrier_action

Represents a split barrier action.

Enumerator
wait 
signal 

Definition at line 292 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

◆ is_channel_enabled()

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

Definition at line 128 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.