DPC++ Runtime
Runtime libraries for oneAPI DPC++
common.hpp
Go to the documentation of this file.
1 //==---------------- common.hpp - DPC++ Explicit SIMD API ----------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 // definitions used in Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
15 #include <cstdint> // for uint* types
16 #include <type_traits>
17 
19 
20 #ifdef __SYCL_DEVICE_ONLY__
21 #define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
22 #define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
23 
24 // Mark a function being nodebug.
25 #define ESIMD_NODEBUG __attribute__((nodebug))
26 // Mark a "ESIMD global": accessible from all functions in current translation
27 // unit, separate copy per subgroup (work-item), mapped to SPIR-V private
28 // storage class.
29 #define ESIMD_PRIVATE \
30  __attribute__((opencl_private)) __attribute__((sycl_explicit_simd))
31 // Bind a ESIMD global variable to a specific register.
32 #define ESIMD_REGISTER(n) __attribute__((register_num(n)))
33 
34 #define __ESIMD_API ESIMD_NODEBUG ESIMD_INLINE
35 
36 #define __ESIMD_UNSUPPORTED_ON_HOST
37 
38 #else // __SYCL_DEVICE_ONLY__
39 #define SYCL_ESIMD_KERNEL
40 #define SYCL_ESIMD_FUNCTION
41 
42 // TODO ESIMD define what this means on Windows host
43 #define ESIMD_NODEBUG
44 // On host device ESIMD global is a thread local static var. This assumes that
45 // each work-item is mapped to a separate OS thread on host device.
46 #define ESIMD_PRIVATE thread_local
47 #define ESIMD_REGISTER(n)
48 
49 #define __ESIMD_API ESIMD_INLINE
50 
51 #define __ESIMD_UNSUPPORTED_ON_HOST \
52  throw sycl::exception(sycl::errc::feature_not_supported, \
53  "This ESIMD feature is not supported on HOST")
54 
55 #endif // __SYCL_DEVICE_ONLY__
56 
57 // Mark a function being noinline
58 #define ESIMD_NOINLINE __attribute__((noinline))
59 // Force a function to be inlined. 'inline' is used to preserve ODR for
60 // functions defined in a header.
61 #define ESIMD_INLINE inline __attribute__((always_inline))
62 
63 // Macros for internal use
64 #define __ESIMD_NS sycl::ext::intel::esimd
65 #define __ESIMD_DNS sycl::ext::intel::esimd::detail
66 #define __ESIMD_EMU_DNS sycl::ext::intel::esimd::emu::detail
67 
68 #define __ESIMD_QUOTE1(m) #m
69 #define __ESIMD_QUOTE(m) __ESIMD_QUOTE1(m)
70 #define __ESIMD_NS_QUOTED __ESIMD_QUOTE(__ESIMD_NS)
71 #define __ESIMD_DEPRECATED(new_api) \
72  __SYCL_DEPRECATED("use " __ESIMD_NS_QUOTED "::" __ESIMD_QUOTE(new_api))
73 
75 
77 namespace __ESIMD_NS {
78 
81 
82 using uchar = unsigned char;
83 using ushort = unsigned short;
84 using uint = unsigned int;
85 
88 struct saturation_on_tag : std::true_type {};
89 
91 struct saturation_off_tag : std::false_type {};
92 
94 static inline constexpr saturation_off_tag saturation_off{};
95 
97 static inline constexpr saturation_on_tag saturation_on{};
98 
100 enum class rgba_channel : uint8_t { R, G, B, A };
101 
105 using SurfaceIndex = unsigned int;
106 
107 namespace detail {
108 template <rgba_channel Ch>
109 static inline constexpr uint8_t ch = 1 << static_cast<int>(Ch);
110 static inline constexpr uint8_t chR = ch<rgba_channel::R>;
111 static inline constexpr uint8_t chG = ch<rgba_channel::G>;
112 static inline constexpr uint8_t chB = ch<rgba_channel::B>;
113 static inline constexpr uint8_t chA = ch<rgba_channel::A>;
114 
115 // Shared Local Memory Binding Table Index (aka surface index).
116 static inline constexpr SurfaceIndex SLM_BTI = 254;
117 static inline constexpr SurfaceIndex INVALID_BTI =
118  static_cast<SurfaceIndex>(-1);
119 } // namespace detail
120 
123 enum class rgba_channel_mask : uint8_t {
124  R = detail::chR,
125  G = detail::chG,
127  B = detail::chB,
131  A = detail::chA,
139 };
140 
142  int Pos = static_cast<int>(Ch);
143  return (static_cast<int>(M) & (1 << Pos)) >> Pos;
144 }
145 
147  return is_channel_enabled(M, rgba_channel::R) +
148  is_channel_enabled(M, rgba_channel::G) +
149  is_channel_enabled(M, rgba_channel::B) +
150  is_channel_enabled(M, rgba_channel::A);
151 }
152 
159 enum class atomic_op : uint8_t {
161  add = 0x0,
163  sub = 0x1,
165  inc = 0x2,
167  dec = 0x3,
169  min = 0x4,
171  max = 0x5,
173  xchg = 0x6,
175  cmpxchg = 0x7,
177  bit_and = 0x8,
179  bit_or = 0x9,
181  bit_xor = 0xa,
183  minsint = 0xb,
185  maxsint = 0xc,
187  fmax = 0x10,
189  fmin = 0x11,
192  fcmpwr = 0x12,
193  fadd = 0x13,
194  fsub = 0x14,
195  load = 0x15,
196  store = 0x16,
199  predec = 0xff,
200 };
201 
203 
204 } // namespace __ESIMD_NS
205 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::intel::esimd::atomic_op
atomic_op
Represents an atomic operation.
Definition: common.hpp:159
cl::sycl::ext::intel::esimd::atomic_op::minsint
@ minsint
Minimum (signed integer): *addr = min(*addr, src0).
cl::sycl::ext::intel::esimd::rgba_channel_mask
rgba_channel_mask
Represents a pixel's channel mask - all possible combinations of enabled channels.
Definition: common.hpp:123
cl::sycl::ext::intel::esimd::atomic_op::xchg
@ xchg
Exchange. *addr == src0;
cl::sycl::ext::intel::esimd::rgba_channel_mask::ABGR
@ ABGR
cl::sycl::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
cl::sycl::ext::intel::esimd::saturation_off_tag
This type tag represents "saturation off" behavior.
Definition: common.hpp:91
cl::sycl::detail::device_global_map::add
void add(const void *DeviceGlobalPtr, const char *UniqueId)
Definition: device_global_map.cpp:16
cl::sycl::ext::intel::esimd::detail::ch
static constexpr uint8_t ch
Definition: common.hpp:109
cl::sycl::ext::intel::esimd::detail::chR
static constexpr uint8_t chR
Definition: common.hpp:110
cl::sycl::ext::intel::esimd::saturation_on
static constexpr saturation_on_tag saturation_on
Type tag object representing "saturation on" behavior.
Definition: common.hpp:97
cl::sycl::ext::intel::esimd::saturation_on_tag
Gen hardware supports applying saturation to results of certain operations.
Definition: common.hpp:88
cl::sycl::ext::intel::esimd::detail::chG
static constexpr uint8_t chG
Definition: common.hpp:111
cl::sycl::ext::intel::esimd::atomic_op::store
@ store
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
cl::sycl::ext::intel::esimd::rgba_channel::R
@ R
cl::sycl::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:21
cl::sycl::ext::intel::esimd::rgba_channel_mask::BGR
@ BGR
cl::sycl::ext::intel::esimd::rgba_channel_mask::AB
@ AB
char
cl::sycl::ext::intel::esimd::rgba_channel_mask::ABG
@ ABG
cl::sycl::ext::intel::esimd::rgba_channel_mask::AR
@ AR
cl::sycl::ext::intel::esimd::rgba_channel_mask::AG
@ AG
cl::sycl::ext::intel::esimd::uchar
unsigned char uchar
Definition: common.hpp:82
cl::sycl::ext::intel::esimd::atomic_op::maxsint
@ maxsint
Maximum (signed integer): *addr = max(*addr, src0).
cl::sycl::ext::intel::esimd::SurfaceIndex
unsigned int SurfaceIndex
Surface index type.
Definition: common.hpp:105
cl::sycl::ext::intel::esimd::get_num_channels_enabled
constexpr int get_num_channels_enabled(rgba_channel_mask M)
Definition: common.hpp:146
cl::sycl::ext::intel::esimd::detail::chA
static constexpr uint8_t chA
Definition: common.hpp:113
cl::sycl::ext::intel::esimd::detail::INVALID_BTI
static constexpr SurfaceIndex INVALID_BTI
Definition: common.hpp:117
cl::sycl::fmax
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmax(T x, T y) __NOEXC
Definition: builtins.hpp:203
defines.hpp
cl::sycl::ext::intel::esimd::rgba_channel_mask::AGR
@ AGR
cl::sycl::fmin
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmin(T x, T y) __NOEXC
Definition: builtins.hpp:216
cl::sycl::ext::intel::esimd::saturation_off
static constexpr saturation_off_tag saturation_off
Type tag object representing "saturation off" behavior.
Definition: common.hpp:94
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::bit_and
std::bit_and< T > bit_and
Definition: functional.hpp:20
cl::sycl::ext::intel::esimd::atomic_op::predec
@ predec
Decrement: *addr = *addr - 1.
cl::sycl::ext::intel::esimd::atomic_op::fadd
@ fadd
cl::sycl::ext::intel::esimd::rgba_channel_mask::ABR
@ ABR
cl::sycl::ext::intel::esimd::rgba_channel_mask::BG
@ BG
cl::sycl::ext::intel::esimd::atomic_op::fcmpwr
@ fcmpwr
Compare and exchange (floating point).
cl::sycl::ext::intel::esimd::rgba_channel_mask::GR
@ GR
cl::sycl::ext::intel::esimd::atomic_op::sub
@ sub
Subtraction: *addr = *addr - src0.
cl::sycl::ext::intel::esimd::atomic_op::cmpxchg
@ cmpxchg
Compare and exchange. if (*addr == src0) *sddr = src1;
cl::sycl::ext::intel::esimd::detail::SLM_BTI
static constexpr SurfaceIndex SLM_BTI
Definition: common.hpp:116
cl::sycl::ext::intel::esimd::atomic_op::load
@ load
cl::sycl::dec
constexpr stream_manipulator dec
Definition: stream.hpp:679
cl::sycl::ext::intel::esimd::rgba_channel_mask::BR
@ BR
cl::sycl::ext::intel::esimd::rgba_channel
rgba_channel
Represents a pixel's channel.
Definition: common.hpp:100
cl::sycl::ext::intel::esimd::ushort
unsigned short ushort
Definition: common.hpp:83
cl::sycl::ext::intel::esimd::detail::chB
static constexpr uint8_t chB
Definition: common.hpp:112
cl::sycl::ext::intel::esimd::rgba_channel::B
@ B
cl::sycl::ext::intel::esimd::rgba_channel::G
@ G
cl::sycl::ext::intel::esimd::atomic_op::inc
@ inc
Increment: *addr = *addr + 1.
cl::sycl::ext::intel::esimd::atomic_op::fsub
@ fsub
cl::sycl::ext::intel::esimd::is_channel_enabled
constexpr int is_channel_enabled(rgba_channel_mask M, rgba_channel Ch)
Definition: common.hpp:141
cl::sycl::ext::intel::esimd::uint
unsigned int uint
Definition: common.hpp:84
cl::sycl::ext::intel::esimd::rgba_channel::A
@ A
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12