Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 5dd32c5

Browse files
AidanBeltonSAidanBeltonS
authored
[SYCL] Update tests for half operator (#1012)
This PR makes some small changes due to the introduction of some missing operators. It changes the namespaces in ext_math.cpp as it was ambiguous as to if the sycl::abs() or sycl::ext::intel::esimd::abs() function should be used. There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats. Depends on: intel/llvm#6061
1 parent 613ede6 commit 5dd32c5

File tree

2 files changed

+23
-21
lines changed

2 files changed

+23
-21
lines changed

‎SYCL/ESIMD/ext_math.cpp‎

Lines changed: 22 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,6 @@
2828

2929
using namespace cl::sycl;
3030
using namespace sycl::ext::intel;
31-
using namespace sycl::ext::intel::esimd;
3231

3332
// --- Data initialization functions
3433

@@ -150,12 +149,13 @@ DEFINE_HOST_BIN_OP(pow, std::pow(X, Y));
150149

151150
#define DEFINE_ESIMD_DEVICE_OP(Op) \
152151
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllVec> { \
153-
simd<T, N> operator()(simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
152+
esimd::simd<T, N> \
153+
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
154154
return esimd::Op<T, N>(X); \
155155
} \
156156
}; \
157157
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllSca> { \
158-
simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
158+
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
159159
return esimd::Op<T, N>(X); \
160160
} \
161161
};
@@ -176,25 +176,26 @@ DEFINE_ESIMD_DEVICE_OP(log2);
176176

177177
#define DEFINE_ESIMD_DEVICE_BIN_OP(Op) \
178178
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllSca> { \
179-
simd<T, N> operator()(T X, \
180-
T Y) const SYCL_ESIMD_FUNCTION { \
179+
esimd::simd<T, N> operator()(T X, T Y) const SYCL_ESIMD_FUNCTION { \
181180
return esimd::Op<T, N>(X, Y); \
182181
} \
183182
}; \
184183
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllVec> { \
185-
simd<T, N> operator()(simd<T, N>X, \
186-
simd<T, N>Y) const SYCL_ESIMD_FUNCTION { \
184+
esimd::simd<T, N> \
185+
operator()(esimd::simd<T, N> X, \
186+
esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
187187
return esimd::Op<T, N>(X, Y); \
188188
} \
189189
}; \
190190
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca1Vec2> { \
191-
simd<T, N> operator()(T X, simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
191+
esimd::simd<T, N> \
192+
operator()(T X, esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
192193
return esimd::Op<T, N>(X, Y); \
193194
} \
194195
}; \
195196
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca2Vec1> { \
196-
simd<T, N> operator()(simd<T, N>X, \
197-
T Y) const SYCL_ESIMD_FUNCTION { \
197+
esimd::simd<T, N> operator()(esimd::simd<T, N> X, \
198+
T Y) const SYCL_ESIMD_FUNCTION { \
198199
return esimd::Op<T, N>(X, Y); \
199200
} \
200201
};
@@ -204,13 +205,14 @@ DEFINE_ESIMD_DEVICE_BIN_OP(pow);
204205

205206
#define DEFINE_SYCL_DEVICE_OP(Op) \
206207
template <class T, int N> struct SYCLf<T, N, MathOp::Op, AllVec> { \
207-
simd<T, N> operator()(simd<T, N>X) const SYCL_ESIMD_FUNCTION { \
208+
esimd::simd<T, N> \
209+
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
208210
/* T must be float for SYCL, so not a template parameter for sycl::Op*/ \
209211
return sycl::Op<N>(X); \
210212
} \
211213
}; \
212214
template <class T, int N> struct SYCLf<T, N, MathOp::Op, AllSca> { \
213-
simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
215+
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
214216
return sycl::Op<N>(X); \
215217
} \
216218
};
@@ -233,14 +235,14 @@ struct UnaryDeviceFunc {
233235

234236
void operator()(id<1> I) const SYCL_ESIMD_KERNEL {
235237
unsigned int Offset = I * N * sizeof(T);
236-
simd<T, N> Vx;
238+
esimd::simd<T, N> Vx;
237239
Vx.copy_from(In, Offset);
238240

239241
if (I.get(0) % 2 == 0) {
240242
for (int J = 0; J < N; J++) {
241243
Kernel<T, N, Op, AllSca> DevF{};
242244
T Val = Vx[J];
243-
simd<T, N> V = DevF(Val); // scalar arg
245+
esimd::simd<T, N> V = DevF(Val); // scalar arg
244246
Vx[J] = V[J];
245247
}
246248
} else {
@@ -264,31 +266,31 @@ struct BinaryDeviceFunc {
264266

265267
void operator()(id<1> I) const SYCL_ESIMD_KERNEL {
266268
unsigned int Offset = I * N * sizeof(T);
267-
simd<T, N> V1(In1, Offset);
268-
simd<T, N> V2(In2, Offset);
269-
simd<T, N> V;
269+
esimd::simd<T, N> V1(In1, Offset);
270+
esimd::simd<T, N> V2(In2, Offset);
271+
esimd::simd<T, N> V;
270272

271273
if (I.get(0) % 2 == 0) {
272274
int Ind = 0;
273275
{
274276
Kernel<T, N, Op, AllSca> DevF{};
275277
T Val2 = V2[Ind];
276-
simd<T, N> Vv = DevF(V1[Ind], Val2); // both arguments are scalar
278+
esimd::simd<T, N> Vv = DevF(V1[Ind], Val2); // both arguments are scalar
277279
V[Ind] = Vv[Ind];
278280
}
279281
Ind++;
280282
{
281283
Kernel<T, N, Op, Sca1Vec2> DevF{};
282284
T Val1 = V1[Ind];
283-
simd<T, N> Vv = DevF(Val1, V2); // scalar, vector
285+
esimd::simd<T, N> Vv = DevF(Val1, V2); // scalar, vector
284286
V[Ind] = Vv[Ind];
285287
}
286288
Ind++;
287289
{
288290
for (int J = Ind; J < N; ++J) {
289291
Kernel<T, N, Op, Sca2Vec1> DevF{};
290292
T Val2 = V2[J];
291-
simd<T, N> Vv = DevF(V1, Val2); // scalar 2nd arg
293+
esimd::simd<T, N> Vv = DevF(V1, Val2); // scalar 2nd arg
292294
V[J] = Vv[J];
293295
}
294296
}

‎SYCL/SubGroup/helper.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -146,7 +146,7 @@ template <typename T> void exit_if_not_equal(T *val, T *ref, const char *name) {
146146
template <> void exit_if_not_equal(half val, half ref, const char *name) {
147147
int16_t cmp_val = reinterpret_cast<int16_t &>(val);
148148
int16_t cmp_ref = reinterpret_cast<int16_t &>(ref);
149-
if (std::abs(cmp_val - cmp_ref) > 1) {
149+
if (std::abs(cmp_val - cmp_ref) > 2) {
150150
std::cout << "Unexpected result for " << name << ": " << (float)val
151151
<< " expected value: " << (float)ref << std::endl;
152152
exit(1);

0 commit comments

Comments
 (0)