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

 cl::sycl::ext::intel::esimd::detail
 
 cl::sycl::ext::intel::experimental::esimd::detail
 

Classes

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

Typedefs

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

Enumerations

enum  cl::sycl::ext::intel::esimd::rgba_channel : uint8_t { cl::sycl::ext::intel::esimd::rgba_channel::R, cl::sycl::ext::intel::esimd::rgba_channel::G, cl::sycl::ext::intel::esimd::rgba_channel::B, cl::sycl::ext::intel::esimd::rgba_channel::A }
 Represents a pixel's channel. More...
 
enum  cl::sycl::ext::intel::esimd::rgba_channel_mask : uint8_t {
  cl::sycl::ext::intel::esimd::rgba_channel_mask::R = detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::G = detail::chG, cl::sycl::ext::intel::esimd::rgba_channel_mask::GR = detail::chG | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::B = detail::chB,
  cl::sycl::ext::intel::esimd::rgba_channel_mask::BR = detail::chB | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::BG = detail::chB | detail::chG, cl::sycl::ext::intel::esimd::rgba_channel_mask::BGR = detail::chB | detail::chG | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::A = detail::chA,
  cl::sycl::ext::intel::esimd::rgba_channel_mask::AR = detail::chA | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::AG = detail::chA | detail::chG, cl::sycl::ext::intel::esimd::rgba_channel_mask::AGR = detail::chA | detail::chG | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::AB = detail::chA | detail::chB,
  cl::sycl::ext::intel::esimd::rgba_channel_mask::ABR = detail::chA | detail::chB | detail::chR, cl::sycl::ext::intel::esimd::rgba_channel_mask::ABG = detail::chA | detail::chB | detail::chG, cl::sycl::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  cl::sycl::ext::intel::esimd::atomic_op : uint8_t {
  cl::sycl::ext::intel::esimd::atomic_op::add = 0x0, cl::sycl::ext::intel::esimd::atomic_op::sub = 0x1, cl::sycl::ext::intel::esimd::atomic_op::inc = 0x2, cl::sycl::ext::intel::esimd::atomic_op::dec = 0x3,
  cl::sycl::ext::intel::esimd::atomic_op::min = 0x4, cl::sycl::ext::intel::esimd::atomic_op::max = 0x5, cl::sycl::ext::intel::esimd::atomic_op::xchg = 0x6, cl::sycl::ext::intel::esimd::atomic_op::cmpxchg = 0x7,
  cl::sycl::ext::intel::esimd::atomic_op::bit_and = 0x8, cl::sycl::ext::intel::esimd::atomic_op::bit_or = 0x9, cl::sycl::ext::intel::esimd::atomic_op::bit_xor = 0xa, cl::sycl::ext::intel::esimd::atomic_op::minsint = 0xb,
  cl::sycl::ext::intel::esimd::atomic_op::maxsint = 0xc, cl::sycl::ext::intel::esimd::atomic_op::fmax = 0x10, cl::sycl::ext::intel::esimd::atomic_op::fmin = 0x11, cl::sycl::ext::intel::esimd::atomic_op::fcmpwr = 0x12,
  cl::sycl::ext::intel::esimd::atomic_op::fadd = 0x13, cl::sycl::ext::intel::esimd::atomic_op::fsub = 0x14, cl::sycl::ext::intel::esimd::atomic_op::load = 0x15, cl::sycl::ext::intel::esimd::atomic_op::store = 0x16,
  cl::sycl::ext::intel::esimd::atomic_op::predec = 0xff
}
 Represents an atomic operation. More...
 
enum  cl::sycl::ext::intel::experimental::esimd::argument_type {
  cl::sycl::ext::intel::experimental::esimd::argument_type::U1 = 0, cl::sycl::ext::intel::experimental::esimd::argument_type::S1 = 1, cl::sycl::ext::intel::experimental::esimd::argument_type::U2 = 2, cl::sycl::ext::intel::experimental::esimd::argument_type::S2 = 3,
  cl::sycl::ext::intel::experimental::esimd::argument_type::U4 = 4, cl::sycl::ext::intel::experimental::esimd::argument_type::S4 = 5, cl::sycl::ext::intel::experimental::esimd::argument_type::U8 = 6, cl::sycl::ext::intel::experimental::esimd::argument_type::S8 = 7,
  cl::sycl::ext::intel::experimental::esimd::argument_type::BF16 = 8, cl::sycl::ext::intel::experimental::esimd::argument_type::FP16 = 9, cl::sycl::ext::intel::experimental::esimd::argument_type::TF32 = 11
}
 
enum  cl::sycl::ext::intel::experimental::esimd::lsc_scope : uint8_t {
  cl::sycl::ext::intel::experimental::esimd::lsc_scope::group = 0, cl::sycl::ext::intel::experimental::esimd::lsc_scope::local = 1, cl::sycl::ext::intel::experimental::esimd::lsc_scope::tile = 2, cl::sycl::ext::intel::experimental::esimd::lsc_scope::gpu = 3,
  cl::sycl::ext::intel::experimental::esimd::lsc_scope::gpus = 4, cl::sycl::ext::intel::experimental::esimd::lsc_scope::system = 5, cl::sycl::ext::intel::experimental::esimd::lsc_scope::sysacq = 6
}
 The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC. More...
 
enum  cl::sycl::ext::intel::experimental::esimd::lsc_fence_op : uint8_t {
  cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::none = 0, cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::evict = 1, cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::invalidate = 2, cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::discard = 3,
  cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::clean = 4, cl::sycl::ext::intel::experimental::esimd::lsc_fence_op::flushl3 = 5
}
 The lsc_fence operation to apply to caches Supported platforms: DG2, PVC. More...
 
enum  cl::sycl::ext::intel::experimental::esimd::lsc_memory_kind : uint8_t { cl::sycl::ext::intel::experimental::esimd::lsc_memory_kind::untyped_global = 0, cl::sycl::ext::intel::experimental::esimd::lsc_memory_kind::untyped_global_low_pri = 1, cl::sycl::ext::intel::experimental::esimd::lsc_memory_kind::typed_global = 2, cl::sycl::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  cl::sycl::ext::intel::experimental::esimd::lsc_data_size : uint8_t {
  cl::sycl::ext::intel::experimental::esimd::lsc_data_size::default_size = 0, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u8 = 1, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u16 = 2, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u32 = 3,
  cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u64 = 4, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u8u32 = 5, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u16u32 = 6, cl::sycl::ext::intel::experimental::esimd::lsc_data_size::u16u32h = 7
}
 Data size or format to read or store. More...
 
enum  cl::sycl::ext::intel::experimental::esimd::cache_hint : uint8_t {
  cl::sycl::ext::intel::experimental::esimd::cache_hint::none = 0, cl::sycl::ext::intel::experimental::esimd::cache_hint::uncached = 1, cl::sycl::ext::intel::experimental::esimd::cache_hint::cached = 2, cl::sycl::ext::intel::experimental::esimd::cache_hint::write_back = 3,
  cl::sycl::ext::intel::experimental::esimd::cache_hint::write_through = 4, cl::sycl::ext::intel::experimental::esimd::cache_hint::streaming = 5, cl::sycl::ext::intel::experimental::esimd::cache_hint::read_invalidate = 6
}
 L1 or L3 cache hint kinds. More...
 
enum  cl::sycl::ext::intel::experimental::esimd::split_barrier_action : uint8_t { cl::sycl::ext::intel::experimental::esimd::split_barrier_action::wait = 0, cl::sycl::ext::intel::experimental::esimd::split_barrier_action::signal = 1 }
 Represents a split barrier action. More...
 

Functions

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

Variables

static constexpr saturation_off_tag cl::sycl::ext::intel::esimd::saturation_off {}
 Type tag object representing "saturation off" behavior. More...
 
static constexpr saturation_on_tag cl::sycl::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.

Typedef Documentation

◆ SurfaceIndex

using cl::sycl::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 105 of file common.hpp.

◆ uchar

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

Definition at line 82 of file common.hpp.

◆ uint

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

Definition at line 84 of file common.hpp.

◆ ushort

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

Definition at line 83 of file common.hpp.

Enumeration Type Documentation

◆ argument_type

Enumerator
U1 
S1 
U2 
S2 
U4 
S4 
U8 
S8 
BF16 
FP16 
TF32 

Definition at line 29 of file common.hpp.

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

min 

Minimum: *addr = min(*addr, src0).

max 

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.

minsint 

Minimum (signed integer): *addr = min(*addr, src0).

maxsint 

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

fmax 

Minimum (floating point): *addr = min(*addr, src0).

fmin 

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

fcmpwr 

Compare and exchange (floating point).

if (*addr == src0) *addr = src1;

fadd 
fsub 
load 
store 
predec 

Decrement: *addr = *addr - 1.

The only operation which returns new value of the destination rather than old.

Definition at line 159 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 338 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 77 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 57 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 69 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 45 of file common.hpp.

◆ rgba_channel

Represents a pixel's channel.

Enumerator

Definition at line 100 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 123 of file common.hpp.

◆ split_barrier_action

Represents a split barrier action.

Enumerator
wait 
signal 

Definition at line 416 of file common.hpp.

Function Documentation

◆ get_num_channels_enabled()

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

Definition at line 146 of file common.hpp.

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

◆ is_channel_enabled()

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

Definition at line 141 of file common.hpp.

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

Variable Documentation

◆ saturation_off

constexpr saturation_off_tag cl::sycl::ext::intel::esimd::saturation_off {}
inlinestaticconstexpr

Type tag object representing "saturation off" behavior.

Definition at line 94 of file common.hpp.

◆ saturation_on

constexpr saturation_on_tag cl::sycl::ext::intel::esimd::saturation_on {}
inlinestaticconstexpr

Type tag object representing "saturation on" behavior.

Definition at line 97 of file common.hpp.