[SYCL] Fix assertion failure in E2E marray test (#14234)

This PR fixes a GPU accuracy bug by upscaling the error-tolerance to a
double type if the GPU supports 64-bit floating point arithmetic.

Author: lbushi25

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

@@ -33,16 +33,28 @@ void test(bool CheckDevice, double delta, FuncTy F, ExpectedTy Expected,
 
   sycl::buffer<bool, 1> SuccessBuf{1};
 
+  sycl::queue q;
+  sycl::device dev = q.get_device();
   // 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) {
+  const bool fp64 = dev.has(sycl::aspect::fp64);
+  q.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, d);
-    });
+    if (fp64) {
+      cgh.single_task([=]() {
+        auto R = F(Args...);
+        static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+        // use double precision error tolerance when fp64 supported
+        Success[0] = equal(R, Expected, delta);
+      });
+    } else {
+      // downscale the error tolerance when fp64 is not supported
+      sycl::detail::get_elem_type_t<ExpectedTy> d(delta);
+      cgh.single_task([=]() {
+        auto R = F(Args...);
+        static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+        Success[0] = equal(R, Expected, d);
+      });
+    }
   });
   assert(sycl::host_accessor{SuccessBuf}[0]);
 }