Skip to content

Commit e8d8b3a

Browse files
kbobrovskbobrovs
authored
[SYCL][ESIMD] Make typenames and constants consistent with SYCL API style. (#3850)
* [SYCL][ESIMD] Make typenames and constants consistent with SYCL API style. Also add some doc comments and a test, improve existing tests. Signed-off-by: kbobrovs <[email protected]>
1 parent 79c460c commit e8d8b3a

File tree

10 files changed

+483
-306
lines changed

10 files changed

+483
-306
lines changed

‎sycl/include/sycl/ext/intel/experimental/esimd.hpp

Lines changed: 1 addition & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -10,16 +10,9 @@
1010

1111
#pragma once
1212

13-
#ifdef __SYCL_DEVICE_ONLY__
14-
#define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
15-
#define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
16-
#else
17-
#define SYCL_ESIMD_KERNEL
18-
#define SYCL_ESIMD_FUNCTION
19-
#endif
20-
2113
/// \defgroup sycl_esimd DPC++ Explicit SIMD API
2214

15+
#include <sycl/ext/intel/experimental/esimd/common.hpp>
2316
#include <sycl/ext/intel/experimental/esimd/math.hpp>
2417
#include <sycl/ext/intel/experimental/esimd/memory.hpp>
2518
#include <sycl/ext/intel/experimental/esimd/simd.hpp>

‎sycl/include/sycl/ext/intel/experimental/esimd/common.hpp

Lines changed: 174 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,14 @@
1414

1515
#include <cstdint> // for uint* types
1616

17+
#ifdef __SYCL_DEVICE_ONLY__
18+
#define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
19+
#define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
20+
#else
21+
#define SYCL_ESIMD_KERNEL
22+
#define SYCL_ESIMD_FUNCTION
23+
#endif
24+
1725
__SYCL_INLINE_NAMESPACE(cl) {
1826
namespace sycl {
1927
namespace ext {
@@ -50,57 +58,162 @@ using uint = unsigned int;
5058
// functions defined in a header.
5159
#define ESIMD_INLINE inline __attribute__((always_inline))
5260

53-
// Enums
54-
// TODO FIXME convert the two enums below to nested enum or class enum to
55-
// remove enum values from the global namespace
56-
enum { GENX_NOSAT = 0, GENX_SAT };
57-
58-
enum ChannelMaskType {
59-
ESIMD_R_ENABLE = 1,
60-
ESIMD_G_ENABLE = 2,
61-
ESIMD_GR_ENABLE = 3,
62-
ESIMD_B_ENABLE = 4,
63-
ESIMD_BR_ENABLE = 5,
64-
ESIMD_BG_ENABLE = 6,
65-
ESIMD_BGR_ENABLE = 7,
66-
ESIMD_A_ENABLE = 8,
67-
ESIMD_AR_ENABLE = 9,
68-
ESIMD_AG_ENABLE = 10,
69-
ESIMD_AGR_ENABLE = 11,
70-
ESIMD_AB_ENABLE = 12,
71-
ESIMD_ABR_ENABLE = 13,
72-
ESIMD_ABG_ENABLE = 14,
73-
ESIMD_ABGR_ENABLE = 15
61+
// Macros for internal use
62+
#define __ESIMD_NS sycl::ext::intel::experimental::esimd
63+
#define __ESIMD_QUOTE1(m) #m
64+
#define __ESIMD_QUOTE(m) __ESIMD_QUOTE1(m)
65+
#define __ESIMD_NS_QUOTED __ESIMD_QUOTE(__ESIMD_NS)
66+
#define __ESIMD_DEPRECATED(new_api) \
67+
__SYCL_DEPRECATED("use " __ESIMD_NS_QUOTED "::" __ESIMD_QUOTE(new_api))
68+
// Defines a deprecated enum value. Use of this value will cause a deprecation
69+
// message printed out by the compiler.
70+
#define __ESIMD_DEPR_ENUM_V(old, new, t) \
71+
old __ESIMD_DEPRECATED(new) = static_cast<t>(new)
72+
73+
/// Gen hardware supports applying saturation to results of some operation.
74+
/// This enum allows to control this behavior.
75+
enum class saturation : uint8_t { off, on };
76+
77+
/// Integer type short-cut to saturation::off.
78+
static inline constexpr uint8_t saturation_off =
79+
static_cast<uint8_t>(saturation::off);
80+
/// Integer type short-cut to saturation::on.
81+
static inline constexpr uint8_t saturation_on =
82+
static_cast<uint8_t>(saturation::on);
83+
84+
enum {
85+
__ESIMD_DEPR_ENUM_V(GENX_NOSAT, saturation::off, uint8_t),
86+
__ESIMD_DEPR_ENUM_V(GENX_SAT, saturation::on, uint8_t)
7487
};
7588

76-
#define NumChannels(Mask) \
77-
((Mask & 1) + ((Mask & 2) >> 1) + ((Mask & 4) >> 2) + ((Mask & 8) >> 3))
78-
79-
#define HasR(Mask) ((Mask & 1) == 1)
80-
#define HasG(Mask) ((Mask & 2) >> 1 == 1)
81-
#define HasB(Mask) ((Mask & 4) >> 2 == 1)
82-
#define HasA(Mask) ((Mask & 8) >> 3 == 1)
83-
84-
enum class EsimdAtomicOpType : uint16_t {
85-
ATOMIC_ADD = 0x0,
86-
ATOMIC_SUB = 0x1,
87-
ATOMIC_INC = 0x2,
88-
ATOMIC_DEC = 0x3,
89-
ATOMIC_MIN = 0x4,
90-
ATOMIC_MAX = 0x5,
91-
ATOMIC_XCHG = 0x6,
92-
ATOMIC_CMPXCHG = 0x7,
93-
ATOMIC_AND = 0x8,
94-
ATOMIC_OR = 0x9,
95-
ATOMIC_XOR = 0xa,
96-
ATOMIC_MINSINT = 0xb,
97-
ATOMIC_MAXSINT = 0xc,
98-
ATOMIC_FMAX = 0x10,
99-
ATOMIC_FMIN = 0x11,
100-
ATOMIC_FCMPWR = 0x12,
101-
ATOMIC_PREDEC = 0xff
89+
/// Represents a pixel's channel.
90+
enum class rgba_channel : uint8_t { R, G, B, A };
91+
92+
namespace detail {
93+
template <rgba_channel Ch>
94+
static inline constexpr uint8_t ch = 1 << static_cast<int>(Ch);
95+
static inline constexpr uint8_t chR = ch<rgba_channel::R>;
96+
static inline constexpr uint8_t chG = ch<rgba_channel::G>;
97+
static inline constexpr uint8_t chB = ch<rgba_channel::B>;
98+
static inline constexpr uint8_t chA = ch<rgba_channel::A>;
99+
} // namespace detail
100+
101+
/// Represents a pixel's channel mask - all possible combinations of enabled
102+
/// channels.
103+
enum class rgba_channel_mask : uint8_t {
104+
R = detail::chR,
105+
G = detail::chG,
106+
GR = detail::chG | detail::chR,
107+
B = detail::chB,
108+
BR = detail::chB | detail::chR,
109+
BG = detail::chB | detail::chG,
110+
BGR = detail::chB | detail::chG | detail::chR,
111+
A = detail::chA,
112+
AR = detail::chA | detail::chR,
113+
AG = detail::chA | detail::chG,
114+
AGR = detail::chA | detail::chG | detail::chR,
115+
AB = detail::chA | detail::chB,
116+
ABR = detail::chA | detail::chB | detail::chR,
117+
ABG = detail::chA | detail::chB | detail::chG,
118+
ABGR = detail::chA | detail::chB | detail::chG | detail::chR,
119+
// For backward compatibility ('ChannelMaskType::ESIMD_R_ENABLE' usage style):
120+
__ESIMD_DEPR_ENUM_V(ESIMD_R_ENABLE, rgba_channel_mask::R, uint8_t),
121+
__ESIMD_DEPR_ENUM_V(ESIMD_G_ENABLE, rgba_channel_mask::G, uint8_t),
122+
__ESIMD_DEPR_ENUM_V(ESIMD_GR_ENABLE, rgba_channel_mask::GR, uint8_t),
123+
__ESIMD_DEPR_ENUM_V(ESIMD_B_ENABLE, rgba_channel_mask::B, uint8_t),
124+
__ESIMD_DEPR_ENUM_V(ESIMD_BR_ENABLE, rgba_channel_mask::BR, uint8_t),
125+
__ESIMD_DEPR_ENUM_V(ESIMD_BG_ENABLE, rgba_channel_mask::BG, uint8_t),
126+
__ESIMD_DEPR_ENUM_V(ESIMD_BGR_ENABLE, rgba_channel_mask::BGR, uint8_t),
127+
__ESIMD_DEPR_ENUM_V(ESIMD_A_ENABLE, rgba_channel_mask::A, uint8_t),
128+
__ESIMD_DEPR_ENUM_V(ESIMD_AR_ENABLE, rgba_channel_mask::AR, uint8_t),
129+
__ESIMD_DEPR_ENUM_V(ESIMD_AG_ENABLE, rgba_channel_mask::AG, uint8_t),
130+
__ESIMD_DEPR_ENUM_V(ESIMD_AGR_ENABLE, rgba_channel_mask::AGR, uint8_t),
131+
__ESIMD_DEPR_ENUM_V(ESIMD_AB_ENABLE, rgba_channel_mask::AB, uint8_t),
132+
__ESIMD_DEPR_ENUM_V(ESIMD_ABR_ENABLE, rgba_channel_mask::ABR, uint8_t),
133+
__ESIMD_DEPR_ENUM_V(ESIMD_ABG_ENABLE, rgba_channel_mask::ABG, uint8_t),
134+
__ESIMD_DEPR_ENUM_V(ESIMD_ABGR_ENABLE, rgba_channel_mask::ABGR, uint8_t)
102135
};
103136

137+
#define __ESIMD_DEPR_CONST(old, new) \
138+
static inline constexpr auto old __ESIMD_DEPRECATED(new) = new
139+
140+
// For backward compatibility ('ESIMD_R_ENABLE' usage style):
141+
__ESIMD_DEPR_CONST(ESIMD_R_ENABLE, rgba_channel_mask::R);
142+
__ESIMD_DEPR_CONST(ESIMD_G_ENABLE, rgba_channel_mask::G);
143+
__ESIMD_DEPR_CONST(ESIMD_GR_ENABLE, rgba_channel_mask::GR);
144+
__ESIMD_DEPR_CONST(ESIMD_B_ENABLE, rgba_channel_mask::B);
145+
__ESIMD_DEPR_CONST(ESIMD_BR_ENABLE, rgba_channel_mask::BR);
146+
__ESIMD_DEPR_CONST(ESIMD_BG_ENABLE, rgba_channel_mask::BG);
147+
__ESIMD_DEPR_CONST(ESIMD_BGR_ENABLE, rgba_channel_mask::BGR);
148+
__ESIMD_DEPR_CONST(ESIMD_A_ENABLE, rgba_channel_mask::A);
149+
__ESIMD_DEPR_CONST(ESIMD_AR_ENABLE, rgba_channel_mask::AR);
150+
__ESIMD_DEPR_CONST(ESIMD_AG_ENABLE, rgba_channel_mask::AG);
151+
__ESIMD_DEPR_CONST(ESIMD_AGR_ENABLE, rgba_channel_mask::AGR);
152+
__ESIMD_DEPR_CONST(ESIMD_AB_ENABLE, rgba_channel_mask::AB);
153+
__ESIMD_DEPR_CONST(ESIMD_ABR_ENABLE, rgba_channel_mask::ABR);
154+
__ESIMD_DEPR_CONST(ESIMD_ABG_ENABLE, rgba_channel_mask::ABG);
155+
__ESIMD_DEPR_CONST(ESIMD_ABGR_ENABLE, rgba_channel_mask::ABGR);
156+
157+
#undef __ESIMD_DEPR_CONST
158+
159+
// For backward compatibility:
160+
using ChannelMaskType = rgba_channel_mask;
161+
162+
constexpr int is_channel_enabled(rgba_channel_mask M, rgba_channel Ch) {
163+
int Pos = static_cast<int>(Ch);
164+
return (static_cast<int>(M) & (1 << Pos)) >> Pos;
165+
}
166+
167+
constexpr int get_num_channels_enabled(rgba_channel_mask M) {
168+
return is_channel_enabled(M, rgba_channel::R) +
169+
is_channel_enabled(M, rgba_channel::G) +
170+
is_channel_enabled(M, rgba_channel::B) +
171+
is_channel_enabled(M, rgba_channel::A);
172+
}
173+
174+
/// Represents an atomic operation.
175+
enum class atomic_op : uint8_t {
176+
add = 0x0,
177+
sub = 0x1,
178+
inc = 0x2,
179+
dec = 0x3,
180+
min = 0x4,
181+
max = 0x5,
182+
xchg = 0x6,
183+
cmpxchg = 0x7,
184+
bit_and = 0x8,
185+
bit_or = 0x9,
186+
bit_xor = 0xa,
187+
minsint = 0xb,
188+
maxsint = 0xc,
189+
fmax = 0x10,
190+
fmin = 0x11,
191+
fcmpwr = 0x12,
192+
predec = 0xff,
193+
// For backward compatibility:
194+
__ESIMD_DEPR_ENUM_V(ATOMIC_ADD, atomic_op::add, uint8_t),
195+
__ESIMD_DEPR_ENUM_V(ATOMIC_SUB, atomic_op::sub, uint8_t),
196+
__ESIMD_DEPR_ENUM_V(ATOMIC_INC, atomic_op::inc, uint8_t),
197+
__ESIMD_DEPR_ENUM_V(ATOMIC_DEC, atomic_op::dec, uint8_t),
198+
__ESIMD_DEPR_ENUM_V(ATOMIC_MIN, atomic_op::min, uint8_t),
199+
__ESIMD_DEPR_ENUM_V(ATOMIC_MAX, atomic_op::max, uint8_t),
200+
__ESIMD_DEPR_ENUM_V(ATOMIC_XCHG, atomic_op::xchg, uint8_t),
201+
__ESIMD_DEPR_ENUM_V(ATOMIC_CMPXCHG, atomic_op::cmpxchg, uint8_t),
202+
__ESIMD_DEPR_ENUM_V(ATOMIC_AND, atomic_op::bit_and, uint8_t),
203+
__ESIMD_DEPR_ENUM_V(ATOMIC_OR, atomic_op::bit_or, uint8_t),
204+
__ESIMD_DEPR_ENUM_V(ATOMIC_XOR, atomic_op::bit_xor, uint8_t),
205+
__ESIMD_DEPR_ENUM_V(ATOMIC_MINSINT, atomic_op::minsint, uint8_t),
206+
__ESIMD_DEPR_ENUM_V(ATOMIC_MAXSINT, atomic_op::maxsint, uint8_t),
207+
__ESIMD_DEPR_ENUM_V(ATOMIC_FMAX, atomic_op::fmax, uint8_t),
208+
__ESIMD_DEPR_ENUM_V(ATOMIC_FMIN, atomic_op::fmin, uint8_t),
209+
__ESIMD_DEPR_ENUM_V(ATOMIC_FCMPWR, atomic_op::fcmpwr, uint8_t),
210+
__ESIMD_DEPR_ENUM_V(ATOMIC_PREDEC, atomic_op::predec, uint8_t)
211+
};
212+
213+
// For backward compatibility:
214+
using EsimdAtomicOpType = atomic_op;
215+
216+
// TODO Cache hints APIs are being reworked.
104217
// L1 or L3 cache hint kinds.
105218
enum class CacheHint : uint8_t {
106219
None = 0,
@@ -111,11 +224,22 @@ enum class CacheHint : uint8_t {
111224
ReadInvalidate = 5
112225
};
113226

114-
enum class EsimdSbarrierType : uint8_t {
115-
WAIT = 0, // split barrier wait
116-
SIGNAL = 1 // split barrier signal
227+
/// Represents a split barrier action.
228+
enum class split_barrier_action : uint8_t {
229+
wait = 0, // split barrier wait
230+
signal = 1, // split barrier signal
231+
// For backward compatibility:
232+
__ESIMD_DEPR_ENUM_V(WAIT, split_barrier_action::wait, uint8_t),
233+
__ESIMD_DEPR_ENUM_V(SIGNAL, split_barrier_action::signal, uint8_t)
117234
};
118235

236+
// For backward compatibility:
237+
using EsimdSbarrierType = split_barrier_action;
238+
239+
#undef __ESIMD_DEPR_ENUM_V
240+
241+
// Since EsimdSbarrierType values are deprecated, these macros will generate
242+
// deprecation message.
119243
#define ESIMD_SBARRIER_WAIT EsimdSbarrierType::WAIT
120244
#define ESIMD_SBARRIER_SIGNAL EsimdSbarrierType::SIGNAL
121245

0 commit comments

Comments
 (0)