[SYCL] Refactor more builtin tests + workaround a bug it uncovered (#12877)

Unfortunately, can't do the same for `sycl::vec` tests because swizzles
can't be captured and a macro seems unavoidable. Also, this includes a
reworked version of #12875 that has
been reverted in #12887.

That also uncovered a bug in our `sycl::length` host implementation when
the project is compiled using GCC. I'm not sure if that's a GCC bug or
yet another instance of UB in `sycl::vec`/`sycl::half` implementation,
but the only option at the moment is to workaround, which I'm doing
here.

Author: aelovikov-intel

sycl/source/builtins/geometric_functions.cpp

@@ -36,11 +36,22 @@ static inline auto dot_host_impl(T0 x, T1 y) {
 }
 EXPORT_SCALAR_AND_VEC_2_4(TWO_ARGS, dot, FP_TYPES)
 
+#if defined(__GNUC__) && !defined(__clang__)
+// GCC miscompiles if using dot (instead of dot_host_impl) *or* if
+// optimizations aren't disabled here. Not sure if a bug in GCC or some UB in
+// sycl::vec/sycl::half (like ansi-alias violations).
+#pragma GCC push_options
+#pragma GCC optimize("O0")
+#endif
 template <typename T> static inline auto length_host_impl(T x) {
-  auto d = dot(x, x);
+  auto d = dot_host_impl(x, x);
   return static_cast<decltype(d)>(std::sqrt(d));
 }
+#if defined(__GNUC__) && !defined(__clang__)
+#pragma GCC pop_options
+#endif
 EXPORT_SCALAR_AND_VEC_2_4(ONE_ARG, length, FP_TYPES)
+
 // fast_length on host is the same as just length.
 template <typename T> static inline auto fast_length_host_impl(T x) {
   return length_host_impl(x);

sycl/test-e2e/Basic/built-ins/helpers.hpp

@@ -1,43 +1,62 @@
 #include <sycl/sycl.hpp>
 
-template <typename T> bool equal(T x, T y) {
+template <typename T, typename D> bool equal(T x, T y, D delta) {
   // Maybe should be C++20's std::equality_comparable.
-  if constexpr (std::is_scalar_v<T>) {
-    return x == y;
+  auto Abs = [](auto x) {
+    if constexpr (std::is_unsigned_v<decltype(x)>)
+      return x;
+    else
+      return std::abs(x);
+  };
+  if constexpr (std::is_scalar_v<T> || std::is_same_v<sycl::half, T>) {
+    return Abs(x - y) <= delta;
   } else {
     for (size_t i = 0; i < x.size(); ++i)
-      if (x[i] != y[i])
+      if (Abs(x[i] - y[i]) > delta)
         return false;
 
     return true;
   }
 }
 
-template <typename FuncTy, typename ExpectedTy,
-          typename... ArgTys>
-void test(bool CheckDevice, FuncTy F, ExpectedTy Expected, ArgTys... Args) {
+template <typename FuncTy, typename ExpectedTy, typename... ArgTys>
+void test(bool CheckDevice, double delta, FuncTy F, ExpectedTy Expected,
+          ArgTys... Args) {
   auto R = F(Args...);
   static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
-  assert(equal(R, Expected));
+  assert(equal(R, Expected, delta));
 
   if (!CheckDevice)
     return;
 
   sycl::buffer<bool, 1> SuccessBuf{1};
+
+  // Make sure we don't use fp64 on devices that don't support it.
+  sycl::detail::get_elem_type_t<ExpectedTy> d(delta);
+
   sycl::queue{}.submit([&](sycl::handler &cgh) {
     sycl::accessor Success{SuccessBuf, cgh};
     cgh.single_task([=]() {
       auto R = F(Args...);
       static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
-      Success[0] = equal(R, Expected);
+      Success[0] = equal(R, Expected, d);
     });
   });
   assert(sycl::host_accessor{SuccessBuf}[0]);
 }
 
 template <typename FuncTy, typename ExpectedTy, typename... ArgTys>
 void test(FuncTy F, ExpectedTy Expected, ArgTys... Args) {
-  test(true /*CheckDevice*/, F, Expected, Args...);
+  test(true /*CheckDevice*/, 0.0 /*delta*/, F, Expected, Args...);
+}
+template <typename FuncTy, typename ExpectedTy,
+          typename... ArgTys>
+void test(bool CheckDevice, FuncTy F, ExpectedTy Expected, ArgTys... Args) {
+  test(CheckDevice, 0.0 /*delta*/, F, Expected, Args...);
+}
+template <typename FuncTy, typename ExpectedTy, typename... ArgTys>
+void test(double delta, FuncTy F, ExpectedTy Expected, ArgTys... Args) {
+  test(true /*CheckDevice*/, delta, F, Expected, Args...);
 }
 
 // MSVC's STL spoils global namespace with math functions, so use explicit

sycl/test-e2e/Basic/built-ins/marray_common.cpp

@@ -8,115 +8,73 @@
 #endif
 #include <cmath>
 
-#include <sycl/sycl.hpp>
-
-#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...)                \
-  {                                                                            \
-    {                                                                          \
-      MARRAY_ELEM_TYPE result[DIM];                                            \
-      {                                                                        \
-        sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM});            \
-        deviceQueue.submit([&](sycl::handler &cgh) {                           \
-          sycl::accessor res_access{b, cgh};                                   \
-          cgh.single_task([=]() {                                              \
-            sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__);       \
-            for (int i = 0; i < DIM; i++)                                      \
-              res_access[i] = res[i];                                          \
-          });                                                                  \
-        });                                                                    \
-      }                                                                        \
-      for (int i = 0; i < DIM; i++)                                            \
-        assert(abs(result[i] - EXPECTED[i]) <= DELTA);                         \
-    }                                                                          \
-  }
-
-#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})
+#include "helpers.hpp"
 
 int main() {
-  sycl::queue deviceQueue;
-  sycl::device dev = deviceQueue.get_device();
+  using namespace sycl;
+
+  marray<float, 2> ma1{1.0f, 2.0f};
+  marray<float, 2> ma2{1.0f, 2.0f};
+  marray<float, 2> ma3{3.0f, 2.0f};
+  marray<double, 2> ma4{1.0, 2.0};
+  marray<float, 3> ma5{M_PI, M_PI, M_PI};
+  marray<double, 3> ma6{M_PI, M_PI, M_PI};
+  marray<half, 3> ma7{M_PI, M_PI, M_PI};
+  marray<float, 2> ma8{0.3f, 0.6f};
+  marray<double, 2> ma9{5.0, 8.0};
+  marray<float, 3> ma10{180, 180, 180};
+  marray<double, 3> ma11{180, 180, 180};
+  marray<half, 3> ma12{180, 180, 180};
+  marray<half, 3> ma13{181, 179, 181};
+  marray<float, 2> ma14{+0.0f, -0.6f};
+  marray<double, 2> ma15{-0.0, 0.6f};
 
-  sycl::marray<float, 2> ma1{1.0f, 2.0f};
-  sycl::marray<float, 2> ma2{1.0f, 2.0f};
-  sycl::marray<float, 2> ma3{3.0f, 2.0f};
-  sycl::marray<double, 2> ma4{1.0, 2.0};
-  sycl::marray<float, 3> ma5{M_PI, M_PI, M_PI};
-  sycl::marray<double, 3> ma6{M_PI, M_PI, M_PI};
-  sycl::marray<sycl::half, 3> ma7{M_PI, M_PI, M_PI};
-  sycl::marray<float, 2> ma8{0.3f, 0.6f};
-  sycl::marray<double, 2> ma9{5.0, 8.0};
-  sycl::marray<float, 3> ma10{180, 180, 180};
-  sycl::marray<double, 3> ma11{180, 180, 180};
-  sycl::marray<sycl::half, 3> ma12{180, 180, 180};
-  sycl::marray<sycl::half, 3> ma13{181, 179, 181};
-  sycl::marray<float, 2> ma14{+0.0f, -0.6f};
-  sycl::marray<double, 2> ma15{-0.0, 0.6f};
+  bool has_fp16 = queue{}.get_device().has(sycl::aspect::fp16);
+  bool has_fp64 = queue{}.get_device().has(sycl::aspect::fp64);
 
-  // sycl::clamp
-  TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3);
-  TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, 1.0f, 3.0f);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::clamp, double, 2, EXPECTED(double, 1.0, 2.0), 0, ma4, 1.0, 3.0);
-  // sycl::degrees
-  TEST(sycl::degrees, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::degrees, double, 3, EXPECTED(double, 180, 180, 180), 0, ma6);
-  if (dev.has(sycl::aspect::fp16))
-    TEST(sycl::degrees, sycl::half, 3, EXPECTED(sycl::half, 180, 180, 180), 0.2,
-         ma7);
-  // sycl::max
-  TEST(sycl::max, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma3);
-  TEST(sycl::max, float, 2, EXPECTED(float, 1.5f, 2.0f), 0, ma1, 1.5f);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::max, double, 2, EXPECTED(double, 1.5, 2.0), 0, ma4, 1.5);
-  // sycl::min
-  TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma3);
-  TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 1.5f), 0, ma1, 1.5f);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::min, double, 2, EXPECTED(double, 1.0, 1.5), 0, ma4, 1.5);
-  // sycl::mix
-  TEST(sycl::mix, float, 2, EXPECTED(float, 1.6f, 2.0f), 0, ma1, ma3, ma8);
-  TEST(sycl::mix, float, 2, EXPECTED(float, 1.4f, 2.0f), 0, ma1, ma3, 0.2);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::mix, double, 2, EXPECTED(double, 3.0, 5.0), 0, ma4, ma9, 0.5);
-  // sycl::radians
-  TEST(sycl::radians, float, 3, EXPECTED(float, M_PI, M_PI, M_PI), 0, ma10);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::radians, double, 3, EXPECTED(double, M_PI, M_PI, M_PI), 0, ma11);
-  if (dev.has(sycl::aspect::fp16))
-    TEST(sycl::radians, sycl::half, 3, EXPECTED(sycl::half, M_PI, M_PI, M_PI),
-         0.002, ma12);
-  // sycl::step
-  TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma1, ma3);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::step, double, 2, EXPECTED(double, 1.0, 1.0), 0, ma4, ma9);
-  if (dev.has(sycl::aspect::fp16))
-    TEST(sycl::step, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), 0,
-         ma12, ma13);
-  TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, 2.5f, ma3);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::step, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, ma9);
-  // sycl::smoothstep
-  TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma8, ma1,
-       ma2);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::smoothstep, double, 2, EXPECTED(double, 1.0, 1.0f), 0.00000001,
-         ma4, ma9, ma9);
-  if (dev.has(sycl::aspect::fp16))
-    TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0),
-         0, ma7, ma12, ma13);
-  TEST(sycl::smoothstep, float, 2, EXPECTED(float, 0.0553936f, 0.0f), 0.0000001,
-       2.5f, 6.0f, ma3);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f,
-         8.0f, ma9);
+  // clamp
+  test(F(clamp), marray<float, 2>{1.0f, 2.0f}, ma1, ma2, ma3);
+  test(F(clamp), marray<float, 2>{1.0f, 2.0f}, ma1, 1.0f, 3.0f);
+  test(has_fp64, F(clamp), marray<double, 2>{1.0, 2.0}, ma4, 1.0, 3.0);
+  // degrees
+  test(F(degrees), marray<float, 3>{180, 180, 180}, ma5);
+  test(has_fp64, F(degrees), marray<double, 3>{180, 180, 180}, ma6);
+  test(has_fp16, 0.2, F(degrees), marray<half, 3>{180, 180, 180}, ma7);
+  // max
+  test(F(max), marray<float, 2>{3.0f, 2.0f}, ma1, ma3);
+  test(F(max), marray<float, 2>{1.5f, 2.0f}, ma1, 1.5f);
+  test(has_fp64, F(max), marray<double, 2>{1.5, 2.0}, ma4, 1.5);
+  // min
+  test(F(min), marray<float, 2>{1.0f, 2.0f}, ma1, ma3);
+  test(F(min), marray<float, 2>{1.0f, 1.5f}, ma1, 1.5f);
+  test(has_fp64, F(min), marray<double, 2>{1.0, 1.5}, ma4, 1.5);
+  // mix
+  test(F(mix), marray<float, 2>{1.6f, 2.0f}, ma1, ma3, ma8);
+  test(F(mix), marray<float, 2>{1.4f, 2.0f}, ma1, ma3, 0.2f);
+  test(has_fp64, F(mix), marray<double, 2>{3.0, 5.0}, ma4, ma9, 0.5);
+  // radians
+  test(F(radians), marray<float, 3>{M_PI, M_PI, M_PI}, ma10);
+  test(has_fp64, F(radians), marray<double, 3>{M_PI, M_PI, M_PI}, ma11);
+  test(has_fp16, 0.002, F(radians), marray<half, 3>{M_PI, M_PI, M_PI}, ma12);
+  // step
+  test(F(step), marray<float, 2>{1.0f, 1.0f}, ma1, ma3);
+  test(has_fp64, F(step), marray<double, 2>{1.0, 1.0}, ma4, ma9);
+  test(has_fp16, F(step), marray<half, 3>{1.0, 0.0, 1.0}, ma12, ma13);
+  test(F(step), marray<float, 2>{1.0f, 0.0f}, 2.5f, ma3);
+  test(has_fp64, F(step), marray<double, 2>{0.0f, 1.0f}, 6.0f, ma9);
+  // smoothstep
+  test(F(smoothstep), marray<float, 2>{1.0f, 1.0f}, ma8, ma1, ma2);
+  test(has_fp64, 0.00000001, F(smoothstep), marray<double, 2>{1.0, 1.0f}, ma4,
+       ma9, ma9);
+  test(has_fp16, F(smoothstep), marray<half, 3>{1.0, 1.0, 1.0}, ma7, ma12,
+       ma13);
+  test(0.0000001, F(smoothstep), marray<float, 2>{0.0553936f, 0.0f}, 2.5f, 6.0f,
+       ma3);
+  test(has_fp64, F(smoothstep), marray<double, 2>{0.0f, 1.0f}, 6.0f, 8.0f, ma9);
   // sign
-  TEST(sycl::sign, float, 2, EXPECTED(float, +0.0f, -1.0f), 0, ma14);
-  if (dev.has(sycl::aspect::fp64))
-    TEST(sycl::sign, double, 2, EXPECTED(double, -0.0, 1.0), 0, ma15);
-  if (dev.has(sycl::aspect::fp16))
-    TEST(sycl::sign, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0), 0,
-         ma12);
+  test(F(sign), marray<float, 2>{+0.0f, -1.0f}, ma14);
+  test(has_fp64, F(sign), marray<double, 2>{-0.0, 1.0}, ma15);
+  test(has_fp16, F(sign), marray<half, 3>{1.0, 1.0, 1.0}, ma12);
 
   return 0;
 }