diff --git a/SYCL/ESIMD/ext_math.cpp b/SYCL/ESIMD/ext_math.cpp index 24b1d53e7b..55a6ee5dfd 100644 --- a/SYCL/ESIMD/ext_math.cpp +++ b/SYCL/ESIMD/ext_math.cpp @@ -28,7 +28,6 @@ using namespace cl::sycl; using namespace sycl::ext::intel; -using namespace sycl::ext::intel::esimd; // --- Data initialization functions @@ -150,12 +149,13 @@ DEFINE_HOST_BIN_OP(pow, std::pow(X, Y)); #define DEFINE_ESIMD_DEVICE_OP(Op) \ template struct ESIMDf { \ - simd operator()(simd X) const SYCL_ESIMD_FUNCTION { \ + esimd::simd \ + operator()(esimd::simd X) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X); \ } \ }; \ template struct ESIMDf { \ - simd operator()(T X) const SYCL_ESIMD_FUNCTION { \ + esimd::simd operator()(T X) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X); \ } \ }; @@ -176,25 +176,26 @@ DEFINE_ESIMD_DEVICE_OP(log2); #define DEFINE_ESIMD_DEVICE_BIN_OP(Op) \ template struct BinESIMDf { \ - simd operator()(T X, \ - T Y) const SYCL_ESIMD_FUNCTION { \ + esimd::simd operator()(T X, T Y) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X, Y); \ } \ }; \ template struct BinESIMDf { \ - simd operator()(simdX, \ - simdY) const SYCL_ESIMD_FUNCTION { \ + esimd::simd \ + operator()(esimd::simd X, \ + esimd::simd Y) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X, Y); \ } \ }; \ template struct BinESIMDf { \ - simd operator()(T X, simd Y) const SYCL_ESIMD_FUNCTION { \ + esimd::simd \ + operator()(T X, esimd::simd Y) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X, Y); \ } \ }; \ template struct BinESIMDf { \ - simd operator()(simdX, \ - T Y) const SYCL_ESIMD_FUNCTION { \ + esimd::simd operator()(esimd::simd X, \ + T Y) const SYCL_ESIMD_FUNCTION { \ return esimd::Op(X, Y); \ } \ }; @@ -204,13 +205,14 @@ DEFINE_ESIMD_DEVICE_BIN_OP(pow); #define DEFINE_SYCL_DEVICE_OP(Op) \ template struct SYCLf { \ - simd operator()(simdX) const SYCL_ESIMD_FUNCTION { \ + esimd::simd \ + operator()(esimd::simd X) const SYCL_ESIMD_FUNCTION { \ /* T must be float for SYCL, so not a template parameter for sycl::Op*/ \ return sycl::Op(X); \ } \ }; \ template struct SYCLf { \ - simd operator()(T X) const SYCL_ESIMD_FUNCTION { \ + esimd::simd operator()(T X) const SYCL_ESIMD_FUNCTION { \ return sycl::Op(X); \ } \ }; @@ -233,14 +235,14 @@ struct UnaryDeviceFunc { void operator()(id<1> I) const SYCL_ESIMD_KERNEL { unsigned int Offset = I * N * sizeof(T); - simd Vx; + esimd::simd Vx; Vx.copy_from(In, Offset); if (I.get(0) % 2 == 0) { for (int J = 0; J < N; J++) { Kernel DevF{}; T Val = Vx[J]; - simd V = DevF(Val); // scalar arg + esimd::simd V = DevF(Val); // scalar arg Vx[J] = V[J]; } } else { @@ -264,23 +266,23 @@ struct BinaryDeviceFunc { void operator()(id<1> I) const SYCL_ESIMD_KERNEL { unsigned int Offset = I * N * sizeof(T); - simd V1(In1, Offset); - simd V2(In2, Offset); - simd V; + esimd::simd V1(In1, Offset); + esimd::simd V2(In2, Offset); + esimd::simd V; if (I.get(0) % 2 == 0) { int Ind = 0; { Kernel DevF{}; T Val2 = V2[Ind]; - simd Vv = DevF(V1[Ind], Val2); // both arguments are scalar + esimd::simd Vv = DevF(V1[Ind], Val2); // both arguments are scalar V[Ind] = Vv[Ind]; } Ind++; { Kernel DevF{}; T Val1 = V1[Ind]; - simd Vv = DevF(Val1, V2); // scalar, vector + esimd::simd Vv = DevF(Val1, V2); // scalar, vector V[Ind] = Vv[Ind]; } Ind++; @@ -288,7 +290,7 @@ struct BinaryDeviceFunc { for (int J = Ind; J < N; ++J) { Kernel DevF{}; T Val2 = V2[J]; - simd Vv = DevF(V1, Val2); // scalar 2nd arg + esimd::simd Vv = DevF(V1, Val2); // scalar 2nd arg V[J] = Vv[J]; } } diff --git a/SYCL/SubGroup/helper.hpp b/SYCL/SubGroup/helper.hpp index 06f239abf0..8d2efb2b43 100644 --- a/SYCL/SubGroup/helper.hpp +++ b/SYCL/SubGroup/helper.hpp @@ -146,7 +146,7 @@ template void exit_if_not_equal(T *val, T *ref, const char *name) { template <> void exit_if_not_equal(half val, half ref, const char *name) { int16_t cmp_val = reinterpret_cast(val); int16_t cmp_ref = reinterpret_cast(ref); - if (std::abs(cmp_val - cmp_ref) > 1) { + if (std::abs(cmp_val - cmp_ref) > 2) { std::cout << "Unexpected result for " << name << ": " << (float)val << " expected value: " << (float)ref << std::endl; exit(1);