Skip to content

Commit 7290092

Browse files
v-klochkovv-klochkov
authored
[ESIMD] Fix unary math functions accepting saturation tag; fix ext_math.cpp test (#10848)
The test ext_math.cpp needed the fix for half type to avoid having only NaNs in inputs. This patch also fixes the comment for saturate function when source and destination type is same FP type. The previous comment wrongly said that saturate() would clamp value to the range [-1.0, 1.0], while ESIMD intrinsic saturate and it is HW version always clamped to [0.0-1.0]. --------- Signed-off-by: Vyacheslav N Klochkov <[email protected]>
1 parent 72b5ecd commit 7290092

File tree

5 files changed

+83
-44
lines changed

5 files changed

+83
-44
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: 59 additions & 42 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,6 @@
99
// RUN: %{build} -fsycl-device-code-split=per_kernel %{mathflags} -o %t.out
1010
// RUN: %{run} %t.out
1111

12-
// FIXME: enable opaque pointers support
13-
// REQUIRES: TEMPORARY_DISABLED
14-
1512
// This test checks extended math operations. Combinations of
1613
// - argument type - half, float
1714
// - math function - sin, cos, ..., div_ieee, pow
@@ -29,6 +26,18 @@
2926
using namespace sycl;
3027
using namespace sycl::ext::intel;
3128

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

3443
// Initialization data for trigonometric functions' input.
@@ -105,22 +114,17 @@ enum class MathOp {
105114

106115
// --- Template functions calculating given math operation on host and device
107116

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

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;
119+
template <class T, int N, MathOp Op, int Args = AllVec> struct ESIMDf;
120+
template <class T, int N, MathOp Op, int Args = AllVec> struct BinESIMDf;
121+
template <class T, int N, MathOp Op, int Args = AllVec> struct SYCLf;
118122

119123
template <class T, MathOp Op> struct HostFunc;
120124

121125
#define DEFINE_HOST_OP(Op, HostOp) \
122126
template <class T> struct HostFunc<T, MathOp::Op> { \
123-
T operator()(T X) { return HostOp; } \
127+
T operator()(T X) { return HOST_SATURATE(HostOp); } \
124128
};
125129

126130
DEFINE_HOST_OP(sin, std::sin(X));
@@ -139,7 +143,7 @@ DEFINE_HOST_OP(log2, std::log2(X));
139143

140144
#define DEFINE_HOST_BIN_OP(Op, HostOp) \
141145
template <class T> struct HostFunc<T, MathOp::Op> { \
142-
T operator()(T X, T Y) { return HostOp; } \
146+
T operator()(T X, T Y) { return HOST_SATURATE(HostOp); } \
143147
};
144148

145149
DEFINE_HOST_BIN_OP(div_ieee, X / Y);
@@ -151,12 +155,12 @@ DEFINE_HOST_BIN_OP(pow, std::pow(X, Y));
151155
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllVec> { \
152156
esimd::simd<T, N> \
153157
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
154-
return esimd::Op<T, N>(X); \
158+
return esimd::Op<T, N>(X, ESIMD_SATURATION_TAG); \
155159
} \
156160
}; \
157161
template <class T, int N> struct ESIMDf<T, N, MathOp::Op, AllSca> { \
158162
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
159-
return esimd::Op<T, N>(X); \
163+
return esimd::Op<T, N>(X, ESIMD_SATURATION_TAG); \
160164
} \
161165
};
162166

@@ -177,26 +181,26 @@ DEFINE_ESIMD_DEVICE_OP(log2);
177181
#define DEFINE_ESIMD_DEVICE_BIN_OP(Op) \
178182
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllSca> { \
179183
esimd::simd<T, N> operator()(T X, T Y) const SYCL_ESIMD_FUNCTION { \
180-
return esimd::Op<T, N>(X, Y); \
184+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
181185
} \
182186
}; \
183187
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, AllVec> { \
184188
esimd::simd<T, N> \
185189
operator()(esimd::simd<T, N> X, \
186190
esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
187-
return esimd::Op<T, N>(X, Y); \
191+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
188192
} \
189193
}; \
190194
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca1Vec2> { \
191195
esimd::simd<T, N> \
192196
operator()(T X, esimd::simd<T, N> Y) const SYCL_ESIMD_FUNCTION { \
193-
return esimd::Op<T, N>(X, Y); \
197+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
194198
} \
195199
}; \
196200
template <class T, int N> struct BinESIMDf<T, N, MathOp::Op, Sca2Vec1> { \
197201
esimd::simd<T, N> operator()(esimd::simd<T, N> X, \
198202
T Y) const SYCL_ESIMD_FUNCTION { \
199-
return esimd::Op<T, N>(X, Y); \
203+
return esimd::Op<T, N>(X, Y, ESIMD_SATURATION_TAG); \
200204
} \
201205
};
202206

@@ -208,12 +212,12 @@ DEFINE_ESIMD_DEVICE_BIN_OP(pow);
208212
esimd::simd<T, N> \
209213
operator()(esimd::simd<T, N> X) const SYCL_ESIMD_FUNCTION { \
210214
/* T must be float for SYCL, so not a template parameter for sycl::Op*/ \
211-
return sycl::Op<N>(X); \
215+
return ESIMD_SATURATE(T, sycl::Op<N>(X)); \
212216
} \
213217
}; \
214218
template <class T, int N> struct SYCLf<T, N, MathOp::Op, AllSca> { \
215219
esimd::simd<T, N> operator()(T X) const SYCL_ESIMD_FUNCTION { \
216-
return sycl::Op<N>(X); \
220+
return ESIMD_SATURATE(T, sycl::Op<N>(X)); \
217221
} \
218222
};
219223

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

313-
constexpr size_t Size = 1024 * 128;
317+
constexpr size_t Size =
318+
std::is_same_v<T, sycl::half> ? (16 * 128) : (1024 * 128);
314319
constexpr bool IsBinOp = (Op == MathOp::div_ieee) || (Op == MathOp::pow);
315320

316321
T *A = new T[Size];
@@ -322,9 +327,9 @@ bool test(queue &Q, const std::string &Name,
322327
Init(A, B, Size);
323328
}
324329
const char *kind =
325-
std::is_same_v<Kernel<T, N, Op, AllVec>, ESIMDf<T, N, Op, AllVec>>
326-
? "ESIMD"
327-
: "SYCL";
330+
std::is_same_v<Kernel<T, N, Op, AllVec>, ESIMDf<T, N, Op, AllVec>>
331+
? "ESIMD"
332+
: "SYCL";
328333
std::cout << " " << Name << " test, kind=" << kind << "...\n";
329334

330335
try {
@@ -343,19 +348,21 @@ bool test(queue &Q, const std::string &Name,
343348
auto PC = BufC.template get_access<access::mode::write>(CGH);
344349
if constexpr (IsBinOp) {
345350
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);
351+
BinaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA, PB,
352+
PC);
348353
CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F);
349354
} else {
350-
UnaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA,
351-
PC);
355+
UnaryDeviceFunc<T, N, Op, Kernel, decltype(PA), decltype(PC)> F(PA, PC);
352356
CGH.parallel_for(nd_range<1>{GlobalRange, LocalRange}, F);
353357
}
354358
});
355359
E.wait();
356360
} catch (sycl::exception &Exc) {
357361
std::cout << " *** ERROR. SYCL exception caught: << " << Exc.what()
358362
<< "\n";
363+
delete[] A;
364+
delete[] B;
365+
delete[] C;
359366
return false;
360367
}
361368

@@ -375,15 +382,19 @@ bool test(queue &Q, const std::string &Name,
375382
}
376383
CheckT Test = C[I];
377384

378-
if (delta == 0.0f) {
379-
delta = sizeof(T) > 2 ? 0.0001 : 0.01;
380-
}
385+
if (delta == 0.0f)
386+
delta = 0.0001;
387+
if constexpr (sizeof(T) <= 2)
388+
delta = delta + delta;
381389

382390
bool BothFinite = std::isfinite(Test) && std::isfinite(Gold);
383391
if (BothFinite && std::abs(Test - Gold) > delta) {
384392
if (++ErrCnt < 10) {
385393
std::cout << " failed at index " << I << ", " << Test
386394
<< " != " << Gold << " (gold)\n";
395+
std::cout << " A = " << (T)A[I] << ", B = " << (T)B[I]
396+
<< ", diff = " << std::abs(Test - Gold)
397+
<< ", max-delta = " << delta << "\n";
387398
}
388399
}
389400
}
@@ -442,10 +453,9 @@ template <class T, int N> bool testESIMDDivIEEE(queue &Q) {
442453

443454
template <class T, int N> bool testESIMDPow(queue &Q) {
444455
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);
456+
std::cout << "--- TESTING ESIMD pow, T=" << typeid(T).name() << ", N = " << N
457+
<< "...\n";
458+
Pass &= test<T, N, MathOp::pow, BinESIMDf>(Q, "pow", InitBin<T>{}, 0.1);
449459
return Pass;
450460
}
451461

@@ -468,9 +478,16 @@ template <class T, int N> bool testSYCL(queue &Q) {
468478

469479
int main(void) {
470480
queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
481+
esimd_test::printTestLabel(Q);
471482
auto Dev = Q.get_device();
472-
std::cout << "Running on " << Dev.get_info<sycl::info::device::name>()
473-
<< "\n";
483+
#ifndef SKIP_NEW_GPU_DRIVER_VERSION_CHECK
484+
if (!esimd_test::isGPUDriverGE(Q, esimd_test::GPUDriverOS::LinuxAndWindows,
485+
"27012", "101.4576")) {
486+
std::cout << "Skipped. The test requires GPU driver 1.3.27012 or newer.\n";
487+
return 0;
488+
}
489+
#endif
490+
474491
bool Pass = true;
475492
#ifdef TEST_IEEE_DIV_REM
476493
Pass &= testESIMDSqrtIEEE<float, 16>(Q);
@@ -479,7 +496,7 @@ int main(void) {
479496
Pass &= testESIMDSqrtIEEE<double, 32>(Q);
480497
Pass &= testESIMDDivIEEE<double, 32>(Q);
481498
}
482-
#else // !TEST_IEEE_DIV_REM
499+
#else // !TEST_IEEE_DIV_REM
483500
Pass &= testESIMD<half, 8>(Q);
484501
Pass &= testESIMD<float, 16>(Q);
485502
Pass &= testESIMD<float, 32>(Q);

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
// The option -fno-slp-vectorize prevents vectorization of code in kernel
2020
// operator() to avoid the extra difficulties in results verification.
2121

22+
#define SKIP_NEW_GPU_DRIVER_VERSION_CHECK 1
2223
#define TEST_FAST_MATH 1
2324

2425
#include "ext_math.cpp"

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,5 +13,6 @@
1313

1414
// This test checks ieee_sqrt() and ieee_sqrt() with float and double types.
1515

16+
#define SKIP_NEW_GPU_DRIVER_VERSION_CHECK 1
1617
#define TEST_IEEE_DIV_REM 1
1718
#include "ext_math.cpp"

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

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
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+
// This test checks extended math operations called with saturation.
13+
// Combinations of
14+
// - argument type - half, float
15+
// - math function - sin, cos, ..., div_ieee, pow
16+
// - SYCL vs ESIMD APIs
17+
18+
#define SATURATION_ON
19+
20+
#include "ext_math.cpp"

0 commit comments

Comments
 (0)