20 #include <type_traits>
24 #ifdef __SYCL_DEVICE_ONLY__
25 #define __ESIMD_UNSUPPORTED_ON_HOST
26 #else // __SYCL_DEVICE_ONLY__
27 #define __ESIMD_UNSUPPORTED_ON_HOST \
28 throw sycl::exception(sycl::errc::feature_not_supported, \
29 "This ESIMD feature is not supported on HOST")
30 #endif // __SYCL_DEVICE_ONLY__
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: