[SYCL] Update part of the tests to make double support optional (#1190)

Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
Co-authored-by: Pavel Chupin <pavel.v.chupin@intel.com>

Author: 3 people

SYCL/AtomicRef/assignment_atomic64.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -10,13 +10,18 @@ using namespace sycl;
 int main() {
   queue q;
 
-  if (!q.get_device().has(aspect::atomic64)) {
+  device dev = q.get_device();
+
+  if (!dev.has(aspect::atomic64)) {
     std::cout << "Skipping test\n";
     return 0;
   }
 
+  const bool DoublesSupported = dev.has(sycl::aspect::fp64);
+
   constexpr int N = 32;
-  assignment_test<double>(q, N);
+  if (DoublesSupported)
+    assignment_test<double>(q, N);
 
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {

SYCL/AtomicRef/assignment_atomic64_generic.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -13,13 +13,18 @@ using namespace sycl;
 int main() {
   queue q;
 
-  if (!q.get_device().has(aspect::atomic64)) {
+  device dev = q.get_device();
+
+  if (!dev.has(aspect::atomic64)) {
     std::cout << "Skipping test\n";
     return 0;
   }
 
+  const bool DoublesSupported = dev.has(sycl::aspect::fp64);
+
   constexpr int N = 32;
-  assignment_generic_test<double>(q, N);
+  if (DoublesSupported)
+    assignment_generic_test<double>(q, N);
 
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {

SYCL/Basic/buffer/buffer.cpp

@@ -510,15 +510,15 @@ int main() {
                                     [](bool *data) { delete[] data; });
     std::shared_ptr<int> int_shrd(new int[size],
                                   [](int *data) { delete[] data; });
-    std::shared_ptr<double> double_shrd(new double[size],
-                                        [](double *data) { delete[] data; });
+    std::shared_ptr<float> float_shrd(new float[size],
+                                      [](float *data) { delete[] data; });
 
     std::vector<bool> bool_vector;
     std::vector<int> int_vector;
-    std::vector<double> double_vector;
+    std::vector<float> float_vector;
     bool_vector.reserve(size);
     int_vector.reserve(size);
-    double_vector.reserve(size);
+    float_vector.reserve(size);
 
     sycl::queue Queue;
     std::mutex m;
@@ -529,40 +529,40 @@ int main() {
       sycl::buffer<int, dims> buf_int_shrd(
           int_shrd, r,
           sycl::property_list{sycl::property::buffer::use_mutex(m)});
-      sycl::buffer<double, dims> buf_double_shrd(
-          double_shrd, r,
+      sycl::buffer<float, dims> buf_float_shrd(
+          float_shrd, r,
           sycl::property_list{sycl::property::buffer::use_mutex(m)});
       m.lock();
       std::fill(bool_shrd.get(), (bool_shrd.get() + size), bool());
       std::fill(int_shrd.get(), (int_shrd.get() + size), int());
-      std::fill(double_shrd.get(), (double_shrd.get() + size), double());
+      std::fill(float_shrd.get(), (float_shrd.get() + size), float());
       m.unlock();
 
       buf_bool_shrd.set_final_data(bool_vector.begin());
       buf_int_shrd.set_final_data(int_vector.begin());
-      buf_double_shrd.set_final_data(double_vector.begin());
+      buf_float_shrd.set_final_data(float_vector.begin());
       buf_bool_shrd.set_write_back(true);
       buf_int_shrd.set_write_back(true);
-      buf_double_shrd.set_write_back(true);
+      buf_float_shrd.set_write_back(true);
 
       Queue.submit([&](sycl::handler &cgh) {
         auto Accessor_bool =
             buf_bool_shrd.get_access<sycl::access::mode::write>(cgh);
         auto Accessor_int =
             buf_int_shrd.get_access<sycl::access::mode::write>(cgh);
-        auto Accessor_double =
-            buf_double_shrd.get_access<sycl::access::mode::write>(cgh);
+        auto Accessor_float =
+            buf_float_shrd.get_access<sycl::access::mode::write>(cgh);
         cgh.parallel_for<class FillBuffer>(r, [=](sycl::id<1> WIid) {
           Accessor_bool[WIid] = true;
           Accessor_int[WIid] = 3;
-          Accessor_double[WIid] = 7.5;
+          Accessor_float[WIid] = 7.5;
         });
       });
     } // Data is copied back
 
     for (size_t i = 0; i < size; i++) {
       if (bool_vector[i] != true || int_vector[i] != 3 ||
-          double_vector[i] != 7.5) {
+          float_vector[i] != 7.5) {
         assert(false && "Data was not copied back");
         return 1;
       }

SYCL/DeviceLib/built-ins/nan.cpp

@@ -1,7 +1,6 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D HALF_IS_SUPPORTED %s -o %t_gpu.out
+// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
-// RUN: %GPU_RUN_PLACEHOLDER %t_gpu.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 #include <iostream>
@@ -59,10 +58,10 @@ int main() {
       }
     }
   });
-#ifdef HALF_IS_SUPPORTED
+
   if (Queue.get_device().has(sycl::aspect::fp16))
     check_nan<unsigned short, s::half>(Queue);
-#endif
+
   check_nan<unsigned int, float>(Queue);
   if (Queue.get_device().has(sycl::aspect::fp64)) {
     check_nan<unsigned long, double>(Queue);

SYCL/ESIMD/api/replicate_smoke.cpp

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
 // The test checks main functionality of the esimd::replicate_vs_w_hs function.
@@ -175,6 +175,7 @@ template <class T> bool test(queue q) {
 int main(int argc, char **argv) {
   queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
   auto dev = q.get_device();
+  const bool doublesSupported = dev.has(sycl::aspect::fp64);
   std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
             << "\n";
   bool passed = true;
@@ -187,7 +188,8 @@ int main(int argc, char **argv) {
   passed &= test<int>(q);
   passed &= test<uint64_t>(q);
   passed &= test<float>(q);
-  passed &= test<double>(q);
+  if (doublesSupported)
+    passed &= test<double>(q);
 
   std::cout << (passed ? "Test passed\n" : "Test FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/api/saturation_smoke.cpp

@@ -9,7 +9,7 @@
 // UNSUPPORTED: cuda || hip
 // TODO: esimd_emulator fails due to unimplemented 'half' type
 // XFAIL: esimd_emulator
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
 // The test checks main functionality of esimd::saturate function.
@@ -184,12 +184,14 @@ int main(int argc, char **argv) {
   auto dev = q.get_device();
   std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
             << "\n";
+  const bool doublesSupported = dev.has(sycl::aspect::fp64);
 
   bool passed = true;
   passed &= test<half, int, FpToInt>(q);
   passed &= test<half, unsigned char, FpToInt>(q);
   passed &= test<float, int, FpToInt>(q);
-  passed &= test<double, short, FpToInt>(q);
+  if (doublesSupported)
+    passed &= test<double, short, FpToInt>(q);
 
   passed &= test<unsigned char, char, UIntToSameOrNarrowAnyInt>(q);
   passed &= test<unsigned short, short, UIntToSameOrNarrowAnyInt>(q);
@@ -206,7 +208,8 @@ int main(int argc, char **argv) {
 
   passed &= test<float, float, FpToFp>(q);
   passed &= test<half, half, FpToFp>(q);
-  passed &= test<double, double, FpToFp>(q);
+  if (doublesSupported)
+    passed &= test<double, double, FpToFp>(q);
 
   std::cout << (passed ? "Test passed\n" : "Test FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/api/simd_view_select_2d_fp.cpp

@@ -24,7 +24,8 @@ int main(int argc, char **argv) {
   bool passed = true;
   passed &= test<half>(q);
   passed &= test<float>(q);
-  passed &= test<double>(q);
+  if (dev.has(sycl::aspect::fp64))
+    passed &= test<double>(q);
 
   std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/ext_math.cpp

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
 // This test checks extended math operations. Combinations of
@@ -474,9 +474,11 @@ int main(void) {
     Pass &= testSYCL<float, 32>(Q);
   }
   Pass &= testESIMDSqrtIEEE<float, 16>(Q);
-  Pass &= testESIMDSqrtIEEE<double, 32>(Q);
+  if (Dev.has(sycl::aspect::fp64)) {
+    Pass &= testESIMDSqrtIEEE<double, 32>(Q);
+    Pass &= testESIMDDivIEEE<double, 32>(Q);
+  }
   Pass &= testESIMDDivIEEE<float, 8>(Q);
-  Pass &= testESIMDDivIEEE<double, 32>(Q);
   Pass &= testESIMDPow<float, 8>(Q);
   Pass &= testESIMDPow<half, 32>(Q);
   std::cout << (Pass ? "Test Passed\n" : "Test FAILED\n");

SYCL/USM/fill.cpp

@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 //
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t1.out
+// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl -fsycl-targets=%sycl_triple  %s -o %t1.out
 // RUN: %CPU_RUN_PLACEHOLDER %t1.out
 // RUN: %GPU_RUN_PLACEHOLDER %t1.out
 // RUN: %ACC_RUN_PLACEHOLDER %t1.out
@@ -124,6 +124,8 @@ int main() {
   auto dev = q.get_device();
   auto ctxt = q.get_context();
 
+  const bool DoublesSupported = dev.has(sycl::aspect::fp64);
+
   test_struct test_obj{4, 42, 424, 4242, 4.2f, 4.242, 4.24242};
 
   if (dev.get_info<info::device::usm_host_allocations>()) {
@@ -133,8 +135,10 @@ int main() {
     runHostTests<long long>(dev, ctxt, q, 4242);
     runHostTests<sycl::half>(dev, ctxt, q, sycl::half(4.2f));
     runHostTests<float>(dev, ctxt, q, 4.242f);
-    runHostTests<double>(dev, ctxt, q, 4.24242);
-    runHostTests<test_struct>(dev, ctxt, q, test_obj);
+    if (DoublesSupported) {
+      runHostTests<double>(dev, ctxt, q, 4.24242);
+      runHostTests<test_struct>(dev, ctxt, q, test_obj);
+    }
   }
 
   if (dev.get_info<info::device::usm_shared_allocations>()) {
@@ -144,8 +148,10 @@ int main() {
     runSharedTests<long long>(dev, ctxt, q, 4242);
     runSharedTests<sycl::half>(dev, ctxt, q, sycl::half(4.2f));
     runSharedTests<float>(dev, ctxt, q, 4.242f);
-    runSharedTests<double>(dev, ctxt, q, 4.24242);
-    runSharedTests<test_struct>(dev, ctxt, q, test_obj);
+    if (DoublesSupported) {
+      runSharedTests<double>(dev, ctxt, q, 4.24242);
+      runSharedTests<test_struct>(dev, ctxt, q, test_obj);
+    }
   }
 
   if (dev.get_info<info::device::usm_device_allocations>()) {
@@ -155,8 +161,10 @@ int main() {
     runDeviceTests<long long>(dev, ctxt, q, 4242);
     runDeviceTests<sycl::half>(dev, ctxt, q, sycl::half(4.2f));
     runDeviceTests<float>(dev, ctxt, q, 4.242f);
-    runDeviceTests<double>(dev, ctxt, q, 4.24242);
-    runDeviceTests<test_struct>(dev, ctxt, q, test_obj);
+    if (DoublesSupported) {
+      runDeviceTests<double>(dev, ctxt, q, 4.24242);
+      runDeviceTests<test_struct>(dev, ctxt, q, test_obj);
+    }
   }
 
   return 0;