Skip to content

Commit 37ddd6c

Browse files
v-klochkovv-klochkov
committed
[ESIMD] Fix the unary math operations accepting saturation tag; fix ext_math.cpp test
The test ext_math.cpp needed the fix for half type to avoid having NaNs in inputs. Signed-off-by: Vyacheslav N Klochkov <[email protected]>
1 parent c043037 commit 37ddd6c

File tree

3 files changed

+66
-36
lines changed

3 files changed

+66
-36
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ namespace ext::intel::esimd {
4646
/// The following conversions are supported:
4747
/// - \c T0 and \c T1 is the same floating-point type (including \c half). In
4848
/// this case the result in the \c i'th lane is:
49-
/// * \c -1 if \c src[i] is less than \c -1
49+
/// * \c 0 if \c src[i] is less than \c 0
5050
/// * \c 1 if \c src[i] is greater than \c 1
5151
/// * src[i] otherwise
5252
///
@@ -352,7 +352,7 @@ ESIMD_NODEBUG
352352
if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
353353
return res; \
354354
else \
355-
return esimd::saturate<T>(res); \
355+
return esimd::saturate<T>(simd<T, N>(res)); \
356356
} \
357357
\
358358
/** Scalar version. */ \

‎sycl/test-e2e/ESIMD/ext_math.cpp

Lines changed: 41 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,18 @@
2929
using namespace sycl;
3030
using namespace sycl::ext::intel;
3131

32+
#ifdef SATURATION_ON
33+
#define ESIMD_SATURATION_TAG \
34+
esimd::saturation_on_tag {}
35+
#define ESIMD_SATURATE(T, x) esimd::saturate<T>(x)
36+
#define HOST_SATURATE(x) std::max(0.0f, std::min((x), 1.0f))
37+
#else
38+
#define ESIMD_SATURATION_TAG \
39+
esimd::saturation_off_tag {}
40+
#define ESIMD_SATURATE(T, x) (x)
41+
#define HOST_SATURATE(x) (x)
42+
#endif
43+
3244
// --- Data initialization functions
3345

3446
// Initialization data for trigonometric functions' input.
@@ -105,22 +117,17 @@ enum class MathOp {
105117

106118
// --- Template functions calculating given math operation on host and device
107119

108-
enum ArgKind {
109-
AllVec,
110-
AllSca,
111-
Sca1Vec2,
112-
Sca2Vec1
113-
};
120+
enum ArgKind { AllVec, AllSca, Sca1Vec2, Sca2Vec1 };
114121

115-
template <class T, int N, MathOp Op, int Args=AllVec> struct ESIMDf;
116-
template <class T, int N, MathOp Op, int Args=AllVec> struct BinESIMDf;
117-
template <class T, int N, MathOp Op, int Args=AllVec> struct SYCLf;
122+
template <class T, int N, MathOp Op, int Args = AllVec> struct ESIMDf;
123+
template <class T, int N, MathOp Op, int Args = AllVec> struct BinESIMDf;
124+
template <class T, int N, MathOp Op, int Args = AllVec> struct SYCLf;
118125

119126
template <class T, MathOp Op> struct HostFunc;
120127

121128
#define DEFINE_HOST_OP(Op, HostOp) \
122129
template <class T> struct HostFunc<T, MathOp::Op> { \
123-
T operator()(T X) { return HostOp; } \
130+
T operator()(T X) { return HOST_SATURATE(HostOp); } \
124131
};
125132

126133
DEFINE_HOST_OP(sin, std::sin(X));
@@ -139,7 +146,7 @@ DEFINE_HOST_OP(log2, std::log2(X));
139146

140147
#define DEFINE_HOST_BIN_OP(Op, HostOp) \
141148
template <class T> struct HostFunc<T, MathOp::Op> { \
142-
T operator()(T X, T Y) { return HostOp; } \
149+
T operator()(T X, T Y) { return HOST_SATURATE(HostOp); } \
143150
};
144151

145152
DEFINE_HOST_BIN_OP(div_ieee, X / Y);
@@ -151,12 +158,12 @@ DEFINE_HOST_BIN_OP(pow, std::pow(X, Y));
151158
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllVec> { \
152159
esimd::simd<T, N> \
153160
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
154-
return esimd::Op<T, N>(X); \
161+
return esimd::Op<T, N>(X, ESIMD_SATURATION_TAG); \
155162
} \
156163
}; \
157164
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllSca> { \
158165
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
159-
return esimd::Op<T, N>(X); \
166+
return esimd::Op<T, N>(X, ESIMD_SATURATION_TAG); \
160167
} \
161168
};
162169

@@ -177,26 +184,26 @@ DEFINE_ESIMD_DEVICE_OP(log2);
177184
#define DEFINE_ESIMD_DEVICE_BIN_OP(Op) \
178185
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllSca> { \
179186
esimd::simd<T, N> operator()(T X, T Y) const SYCL_ESIMD_FUNCTION { \
180-
return esimd::Op<T, N>(X, Y); \
187+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
181188
} \
182189
}; \
183190
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllVec> { \
184191
esimd::simd<T, N> \
185192
operator()(esimd::simd<T, N> X, \
186193
esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
187-
return esimd::Op<T, N>(X, Y); \
194+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
188195
} \
189196
}; \
190197
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca1Vec2> { \
191198
esimd::simd<T, N> \
192199
operator()(T X, esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
193-
return esimd::Op<T, N>(X, Y); \
200+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
194201
} \
195202
}; \
196203
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca2Vec1> { \
197204
esimd::simd<T, N> operator()(esimd::simd<T, N> X, \
198205
T Y) const SYCL_ESIMD_FUNCTION { \
199-
return esimd::Op<T, N>(X, Y); \
206+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
200207
} \
201208
};
202209

@@ -208,12 +215,12 @@ DEFINE_ESIMD_DEVICE_BIN_OP(pow);
208215
esimd::simd<T, N> \
209216
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
210217
/* T must be float for SYCL, so not a template parameter for sycl::Op*/ \
211-
return sycl::Op<N>(X); \
218+
return ESIMD_SATURATE(T, sycl::Op<N>(X)); \
212219
} \
213220
}; \
214221
template <class T, int N> struct SYCLf<T, N, MathOp::Op, AllSca> { \
215222
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
216-
return sycl::Op<N>(X); \
223+
return ESIMD_SATURATE(T, sycl::Op<N>(X)); \
217224
} \
218225
};
219226

@@ -307,10 +314,11 @@ struct BinaryDeviceFunc {
307314
template <class T, int N, MathOp Op,
308315
template <class, int, MathOp, int> class Kernel,
309316
typename InitF = InitNarrow<T>>
310-
bool test(queue &Q, const std::string &Name,
311-
InitF Init = InitNarrow<T>{}, float delta = 0.0f) {
317+
bool test(queue &Q, const std::string &Name, InitF Init = InitNarrow<T>{},
318+
float delta = 0.0f) {
312319

313-
constexpr size_t Size = 1024 * 128;
320+
constexpr size_t Size =
321+
std::is_same_v<T, sycl::half> ? (16 * 128) : (1024 * 128);
314322
constexpr bool IsBinOp = (Op == MathOp::div_ieee) || (Op == MathOp::pow);
315323

316324
T *A = new T[Size];
@@ -322,9 +330,9 @@ bool test(queue &Q, const std::string &Name,
322330
Init(A, B, Size);
323331
}
324332
const char *kind =
325-
std::is_same_v<Kernel<T, N, Op, AllVec>, ESIMDf<T, N, Op, AllVec>>
326-
? "ESIMD"
327-
: "SYCL";
333+
std::is_same_v<Kernel<T, N, Op, AllVec>, ESIMDf<T, N, Op, AllVec>>
334+
? "ESIMD"
335+
: "SYCL";
328336
std::cout << " " << Name << " test, kind=" << kind << "...\n";
329337

330338
try {
@@ -343,12 +351,11 @@ bool test(queue &Q, const std::string &Name,
343351
auto PC = BufC.template get_access<access::mode::write>(CGH);
344352
if constexpr (IsBinOp) {
345353
auto PB = BufB.template get_access<access::mode::read>(CGH);
346-
BinaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(
347-
PA, PB, PC);
354+
BinaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA, PB,
355+
PC);
348356
CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F);
349357
} else {
350-
UnaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA,
351-
PC);
358+
UnaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA, PC);
352359
CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F);
353360
}
354361
});
@@ -384,6 +391,7 @@ bool test(queue &Q, const std::string &Name,
384391
if (++ErrCnt < 10) {
385392
std::cout << " failed at index " << I << ", " << Test
386393
<< " != " << Gold << " (gold)\n";
394+
std::cout << "A = " << (T)A[I] << ", B = " << (T)B[I] << "\n";
387395
}
388396
}
389397
}
@@ -442,10 +450,9 @@ template <class T, int N> bool testESIMDDivIEEE(queue &Q) {
442450

443451
template <class T, int N> bool testESIMDPow(queue &Q) {
444452
bool Pass = true;
445-
std::cout << "--- TESTING ESIMD pow, T=" << typeid(T).name()
446-
<< ", N = " << N << "...\n";
447-
Pass &= test<T, N, MathOp::pow, BinESIMDf>(
448-
Q, "pow", InitBin<T>{}, 0.1);
453+
std::cout << "--- TESTING ESIMD pow, T=" << typeid(T).name() << ", N = " << N
454+
<< "...\n";
455+
Pass &= test<T, N, MathOp::pow, BinESIMDf>(Q, "pow", InitBin<T>{}, 0.1);
449456
return Pass;
450457
}
451458

@@ -479,7 +486,7 @@ int main(void) {
479486
Pass &= testESIMDSqrtIEEE<double, 32>(Q);
480487
Pass &= testESIMDDivIEEE<double, 32>(Q);
481488
}
482-
#else // !TEST_IEEE_DIV_REM
489+
#else // !TEST_IEEE_DIV_REM
483490
Pass &= testESIMD<half, 8>(Q);
484491
Pass &= testESIMD<float, 16>(Q);
485492
Pass &= testESIMD<float, 32>(Q);

‎sycl/test-e2e/ESIMD/ext_math_saturate.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
//==----------- ext_math_saturate.cpp - DPC++ ESIMD extended math test ----==//
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+
// DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
9+
// RUN: %{build} -fsycl-device-code-split=per_kernel %{mathflags} -o %t.out
10+
// RUN: %{run} %t.out
11+
12+
// FIXME: enable opaque pointers support
13+
// REQUIRES: TEMPORARY_DISABLED
14+
15+
// This test checks extended math operations called with saturation.
16+
// Combinations of
17+
// - argument type - half, float
18+
// - math function - sin, cos, ..., div_ieee, pow
19+
// - SYCL vs ESIMD APIs
20+
21+
#define SATURATION_ON
22+
23+
#include "ext_math.cpp"

0 commit comments

Comments
 (0)