20 #include <type_traits>
24 #ifdef __SYCL_DEVICE_ONLY__
25 #define __ESIMD_UNSUPPORTED_ON_HOST
27 #define __ESIMD_UNSUPPORTED_ON_HOST \
28 throw sycl::exception(sycl::errc::feature_not_supported, \
29 "This ESIMD feature is not supported on HOST")
36 namespace ext::intel::esimd {
70 static constexpr
bool value =
71 std::is_same_v<T, __ESIMD_NS::saturation_on_tag> ||
72 std::is_same_v<T, __ESIMD_NS::saturation_off_tag>;
80 return (n & (n - 1)) == 0;
86 ESIMD_INLINE constexpr
bool isPowerOf2(
unsigned int n,
unsigned int limit) {
87 return (n & (n - 1)) == 0 && n <= limit;
90 template <rgba_channel Ch>
91 static inline constexpr uint8_t
ch = 1 <<
static_cast<int>(Ch);
92 static inline constexpr uint8_t
chR = ch<rgba_channel::R>;
93 static inline constexpr uint8_t
chG = ch<rgba_channel::G>;
94 static inline constexpr uint8_t
chB = ch<rgba_channel::B>;
95 static inline constexpr uint8_t
chA = ch<rgba_channel::A>;
124 int Pos =
static_cast<int>(Ch);
125 return (
static_cast<int>(M) & (1 << Pos)) >> Pos;
135 #define __ESIMD_USM_DWORD_ATOMIC_TO_LSC \
136 " is supported only on ACM, PVC. USM-based atomic will be auto-converted " \
193 #undef __ESIMD_USM_DWORD_TO_LSC_MSG
198 template <__ESIMD_NS::native::lsc::atomic_op Op> constexpr
int get_num_args() {
199 if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::inc ||
201 Op == __ESIMD_NS::native::lsc::atomic_op::load) {
203 }
else if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::store ||
205 Op == __ESIMD_NS::native::lsc::atomic_op::sub ||
206 Op == __ESIMD_NS::native::lsc::atomic_op::smin ||
207 Op == __ESIMD_NS::native::lsc::atomic_op::smax ||
208 Op == __ESIMD_NS::native::lsc::atomic_op::umin ||
209 Op == __ESIMD_NS::native::lsc::atomic_op::umax ||
210 Op == __ESIMD_NS::native::lsc::atomic_op::fadd ||
211 Op == __ESIMD_NS::native::lsc::atomic_op::fsub ||
218 }
else if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::cmpxchg ||
219 Op == __ESIMD_NS::native::lsc::atomic_op::fcmpxchg) {
228 case __ESIMD_NS::atomic_op::xchg:
229 case __ESIMD_NS::atomic_op::predec:
236 template <__ESIMD_NS::atomic_op Op>
241 case __ESIMD_NS::atomic_op::sub:
242 return __ESIMD_NS::native::lsc::atomic_op::sub;
243 case __ESIMD_NS::atomic_op::inc:
244 return __ESIMD_NS::native::lsc::atomic_op::inc;
248 return __ESIMD_NS::native::lsc::atomic_op::umin;
250 return __ESIMD_NS::native::lsc::atomic_op::umax;
251 case __ESIMD_NS::atomic_op::cmpxchg:
252 return __ESIMD_NS::native::lsc::atomic_op::cmpxchg;
259 case __ESIMD_NS::atomic_op::minsint:
260 return __ESIMD_NS::native::lsc::atomic_op::smin;
261 case __ESIMD_NS::atomic_op::maxsint:
262 return __ESIMD_NS::native::lsc::atomic_op::smax;
267 case __ESIMD_NS::atomic_op::fcmpwr:
268 return __ESIMD_NS::native::lsc::atomic_op::fcmpxchg;
269 case __ESIMD_NS::atomic_op::fadd:
270 return __ESIMD_NS::native::lsc::atomic_op::fadd;
271 case __ESIMD_NS::atomic_op::fsub:
272 return __ESIMD_NS::native::lsc::atomic_op::fsub;
273 case __ESIMD_NS::atomic_op::load:
274 return __ESIMD_NS::native::lsc::atomic_op::load;
275 case __ESIMD_NS::atomic_op::store:
276 return __ESIMD_NS::native::lsc::atomic_op::store;
278 static_assert(has_lsc_equivalent<Op>() &&
"Unsupported LSC atomic op");
282 template <__ESIMD_NS::native::lsc::atomic_op Op>
287 case __ESIMD_NS::native::lsc::atomic_op::sub:
288 return __ESIMD_NS::atomic_op::sub;
289 case __ESIMD_NS::native::lsc::atomic_op::inc:
290 return __ESIMD_NS::atomic_op::inc;
293 case __ESIMD_NS::native::lsc::atomic_op::umin:
295 case __ESIMD_NS::native::lsc::atomic_op::umax:
297 case __ESIMD_NS::native::lsc::atomic_op::cmpxchg:
298 return __ESIMD_NS::atomic_op::cmpxchg;
305 case __ESIMD_NS::native::lsc::atomic_op::smin:
306 return __ESIMD_NS::atomic_op::minsint;
307 case __ESIMD_NS::native::lsc::atomic_op::smax:
308 return __ESIMD_NS::atomic_op::maxsint;
313 case __ESIMD_NS::native::lsc::atomic_op::fcmpxchg:
314 return __ESIMD_NS::atomic_op::fcmpwr;
315 case __ESIMD_NS::native::lsc::atomic_op::fadd:
316 return __ESIMD_NS::atomic_op::fadd;
317 case __ESIMD_NS::native::lsc::atomic_op::fsub:
318 return __ESIMD_NS::atomic_op::fsub;
319 case __ESIMD_NS::native::lsc::atomic_op::load:
320 return __ESIMD_NS::atomic_op::load;
321 case __ESIMD_NS::native::lsc::atomic_op::store:
322 return __ESIMD_NS::atomic_op::store;
326 template <__ESIMD_NS::atomic_op Op> constexpr
int get_num_args() {
327 if constexpr (has_lsc_equivalent<Op>()) {
328 return get_num_args<to_lsc_atomic_op<Op>()>();
331 case __ESIMD_NS::atomic_op::xchg:
332 case __ESIMD_NS::atomic_op::predec:
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL_DEPRECATED(message)
#define __ESIMD_USM_DWORD_ATOMIC_TO_LSC
constexpr int is_channel_enabled(rgba_channel_mask M, rgba_channel Ch)
rgba_channel
Represents a pixel's channel.
rgba_channel_mask
Represents a pixel's channel mask - all possible combinations of enabled channels.
unsigned int SurfaceIndex
Surface index type.
static constexpr saturation_off_tag saturation_off
Type tag object representing "saturation off" behavior.
constexpr int get_num_channels_enabled(rgba_channel_mask M)
static constexpr saturation_on_tag saturation_on
Type tag object representing "saturation on" behavior.
atomic_op
Represents an atomic operation.
@ umin
Minimum: *addr = min(*addr, src0).
@ smin
Minimum (signed integer): *addr = min(*addr, src0).
@ cmpxchg
Compare and exchange. if (*addr == src0) *sddr = src1;
@ umax
Maximum: *addr = max(*addr, src0).
@ sub
Subtraction: *addr = *addr - src0.
@ xchg
Exchange. *addr == src0;
@ predec
Decrement: *addr = *addr - 1.
@ smax
Maximum (signed integer): *addr = max(*addr, src0).
@ inc
Increment: *addr = *addr + 1.
@ fcmpxchg
Compare and exchange (floating point).
atomic_op
LSC atomic operation codes.
void add(const void *DeviceGlobalPtr, const char *UniqueId)
constexpr sycl::ext::intel::esimd::atomic_op to_atomic_op()
constexpr bool has_lsc_equivalent()
constexpr ESIMD_INLINE bool isPowerOf2(unsigned int n, unsigned int limit)
Check at compile time if given 32 bit positive integer is both:
static constexpr SurfaceIndex SLM_BTI
constexpr sycl::ext::intel::esimd::native::lsc::atomic_op to_lsc_atomic_op()
constexpr bool is_saturation_tag_v
static constexpr uint8_t chG
static constexpr uint8_t chR
static constexpr uint8_t chB
static constexpr uint8_t ch
static constexpr uint8_t chA
constexpr int get_num_args()
static constexpr SurfaceIndex INVALID_BTI
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmax(T x, T y)
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmin(T x, T y)
std::bit_xor< T > bit_xor
std::bit_and< T > bit_and
std::bit_and< T > bit_and
std::bit_xor< T > bit_xor
constexpr stream_manipulator dec
---— Error handling, matching OpenCL plugin semantics.
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
This type tag represents "saturation off" behavior.
Gen hardware supports applying saturation to results of certain operations.