[SYCL] Split 'double type' test into two tests, 1) <test>.cpp; 2) <test>_aspect-fp64.cpp for llvm-test-suite

SYCL/AtomicRef/assignment_atomic64.cpp

@@ -16,7 +16,9 @@ int main() {
   }
 
   constexpr int N = 32;
+#ifdef ENABLE_FP64
   assignment_test<double>(q, N);
+#endif
 
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {

SYCL/AtomicRef/assignment_atomic64_aspect_fp64.cpp

@@ -0,0 +1,16 @@
+// Enable FP64 part of <assignment_atomic64.cpp>. To be removed once DPC++
+// supports optional device features and the code could be enabled
+// unconditionally without causing failures in speculative compilation
+// of the kernels.
+//
+// REQUIRES: aspect-fp64
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DENABLE_FP64 %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "assignment_atomic64.cpp"

SYCL/AtomicRef/assignment_atomic64_generic.cpp

@@ -19,8 +19,9 @@ int main() {
   }
 
   constexpr int N = 32;
+#ifdef ENABLE_FP64
   assignment_generic_test<double>(q, N);
-
+#endif
   // Include long tests if they are 64 bits wide
   if constexpr (sizeof(long) == 8) {
     assignment_generic_test<long>(q, N);
@@ -37,6 +38,5 @@ int main() {
   if constexpr (sizeof(char *) == 8) {
     assignment_generic_test<char *>(q, N);
   }
-
   std::cout << "Test passed." << std::endl;
 }

SYCL/AtomicRef/assignment_atomic64_generic_aspect_fp64.cpp

@@ -0,0 +1,16 @@
+// Enable FP64 part of <assignment_atomic64_generic.cpp>. To be removed once
+// DPC++ supports optional device features and the code could be enabled
+// unconditionally without causing failures in speculative compilation
+// of the kernels.
+//
+// REQUIRES: aspect-fp64
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DENABLE_FP64 %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// CUDA backend has had no support for the generic address space yet
+// XFAIL: cuda || hip
+
+#include "assignment_atomic64_generic.cpp"

SYCL/Basic/buffer/buffer.cpp

@@ -14,10 +14,10 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <sycl/sycl.hpp>
+
 #include <cassert>
-#include <iostream>
 #include <memory>
-#include <sycl/sycl.hpp>
 
 using namespace sycl;
 
@@ -513,15 +513,22 @@ int main() {
                                     [](bool *data) { delete[] data; });
     std::shared_ptr<int> int_shrd(new int[size],
                                   [](int *data) { delete[] data; });
+#ifdef ENABLE_FP64
     std::shared_ptr<double> double_shrd(new double[size],
                                         [](double *data) { delete[] data; });
+#endif
 
     std::vector<bool> bool_vector;
     std::vector<int> int_vector;
+#ifdef ENABLE_FP64
     std::vector<double> double_vector;
+#endif
+
     bool_vector.reserve(size);
     int_vector.reserve(size);
+#ifdef ENABLE_FP64
     double_vector.reserve(size);
+#endif
 
     sycl::queue Queue;
     std::mutex m;
@@ -532,43 +539,58 @@ int main() {
       sycl::buffer<int, dims> buf_int_shrd(
           int_shrd, r,
           sycl::property_list{sycl::property::buffer::use_mutex(m)});
+#ifdef ENABLE_FP64
       sycl::buffer<double, dims> buf_double_shrd(
           double_shrd, r,
           sycl::property_list{sycl::property::buffer::use_mutex(m)});
+#endif
       m.lock();
       std::fill(bool_shrd.get(), (bool_shrd.get() + size), bool());
       std::fill(int_shrd.get(), (int_shrd.get() + size), int());
+#ifdef ENABLE_FP64
       std::fill(double_shrd.get(), (double_shrd.get() + size), double());
+#endif
       m.unlock();
 
       buf_bool_shrd.set_final_data(bool_vector.begin());
       buf_int_shrd.set_final_data(int_vector.begin());
+#ifdef ENABLE_FP64
       buf_double_shrd.set_final_data(double_vector.begin());
+#endif
+
       buf_bool_shrd.set_write_back(true);
       buf_int_shrd.set_write_back(true);
+#ifdef ENABLE_FP64
       buf_double_shrd.set_write_back(true);
+#endif
 
       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);
+#ifdef ENABLE_FP64
         auto Accessor_double =
             buf_double_shrd.get_access<sycl::access::mode::write>(cgh);
+#endif
         cgh.parallel_for<class FillBuffer>(r, [=](sycl::id<1> WIid) {
           Accessor_bool[WIid] = true;
           Accessor_int[WIid] = 3;
+#ifdef ENABLE_FP64
           Accessor_double[WIid] = 7.5;
+#endif
         });
       });
     } // 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) {
-        assert(false && "Data was not copied back");
-        return 1;
-      }
+      bool Passed = true;
+      Passed &= (bool_vector[i] == true);
+      Passed &= (int_vector[i] == 3);
+#ifdef ENABLE_FP64
+      Passed &= (double_vector[i] == 7.5);
+#endif
+      assert(Passed && "Data was not copied back");
     }
   }
 

SYCL/Basic/buffer/buffer_aspect_fp64.cpp

@@ -0,0 +1,23 @@
+// Enable FP64 part of <buffer.cpp>. To be removed once DPC++
+// supports optional device features and the code could be enabled
+// unconditionally without causing failures in speculative compilation
+// of the kernels.
+//
+// REQUIRES: aspect-fp64
+// RUN: %clangxx %cxx_std_optionc++17 -DENABLE_FP64 %s -o %t1.out %sycl_options
+// RUN: %HOST_RUN_PLACEHOLDER %t1.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t2.out
+// RUN: %HOST_RUN_PLACEHOLDER %t2.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2.out
+
+//==------------------- buffer.cpp - SYCL buffer basic test ----------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "buffer.cpp"

SYCL/DeviceLib/built-ins/nan.cpp

@@ -41,13 +41,14 @@ template <typename T, typename R> void check_nan(s::queue &Queue) {
 int main() {
   test_nan_call<s::ushort, s::half>();
   test_nan_call<s::uint, float>();
-  test_nan_call<s::ulong, double>();
-  test_nan_call<s::ulonglong, double>();
   test_nan_call<s::ushort2, s::half2>();
   test_nan_call<s::uint2, s::float2>();
+#ifdef ENABLE_FP64
+  test_nan_call<s::ulong, double>();
+  test_nan_call<s::ulonglong, double>();
   test_nan_call<s::ulong2, s::double2>();
   test_nan_call<s::ulonglong2, s::double2>();
-
+#endif
   s::queue Queue([](sycl::exception_list ExceptionList) {
     for (std::exception_ptr ExceptionPtr : ExceptionList) {
       try {
@@ -63,10 +64,13 @@ int main() {
   if (Queue.get_device().has(sycl::aspect::fp16))
     check_nan<unsigned short, s::half>(Queue);
 #endif
+
   check_nan<unsigned int, float>(Queue);
+#ifdef ENABLE_FP64
   if (Queue.get_device().has(sycl::aspect::fp64)) {
     check_nan<unsigned long, double>(Queue);
     check_nan<unsigned long long, double>(Queue);
   }
+#endif
   return 0;
 }

SYCL/DeviceLib/built-ins/nan_aspect_fp64.cpp

@@ -0,0 +1,13 @@
+// Enable FP64 part of <nan.cpp>. To be removed once DPC++
+// supports optional device features and the code could be enabled
+// unconditionally without causing failures in speculative compilation
+// of the kernels.
+//
+// REQUIRES: aspect-fp64
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DENABLE_FP64 %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_gpu.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include "nan.cpp"

SYCL/ESIMD/aot_mixed.cpp

@@ -8,9 +8,9 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
 // UNSUPPORTED: esimd_emulator
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" -o %t.sycl.out -DENABLE_SYCL=0 %s
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen12lp" -o %t.sycl.out -DENABLE_SYCL=0 %s
 // RUN: %GPU_RUN_PLACEHOLDER %t.sycl.out
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" -o %t.out %s
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device gen12lp" -o %t.out %s
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
 // This test checks the following ESIMD ahead-of-time compilation scenarios:

SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp

@@ -265,12 +265,17 @@ int main(void) {
   auto arith_ops = esimd_test::ArithBinaryOpsNoDiv;
   passed &= test<unsigned char, int, 1, BinOp, VSf, IDf>(arith_ops, q);
   passed &= test<char, float, 7, BinOp, VEf, IDf>(arith_ops, q, 0.000001f);
+#ifdef ENABLE_FP64
   passed &= test<short, double, 7, BinOp, VEf, IDf>(arith_ops, q, 1e-15);
+#endif
   passed &= test<float, float, 32, BinOp, VEf, IDf>(arith_ops, q, 0.000001f);
   passed &= test<half, char, 1, BinOp, verify_n, IDf>(arith_ops, q, 1);
   passed &= test<half, unsigned int, 32, BinOp, VSf, IDf>(arith_ops, q, 1);
-  passed &= test<double, half, 7, BinOp, VSf, IDf>(arith_ops, q);
   passed &= test<short, uint64_t, 7, BinOp, VSf, IDf>(arith_ops, q);
+#ifdef ENABLE_FP64
+  passed &= test<double, half, 7, BinOp, VSf, IDf>(arith_ops, q);
+  passed &= test<short, double, 7, BinOp, VSf, IDf>(arith_ops, q);
+#endif
 #ifdef USE_BF16
   passed &= test<bfloat16, int, 8, BinOp, VSf, IDf>(arith_ops, q);
   passed &= test<half, bfloat16, 7, BinOp, VEfa, IDf>(arith_ops, q, 0.03);
@@ -326,12 +331,17 @@ int main(void) {
   auto cmp_ops = esimd_test::CmpOps;
   passed &= test<unsigned char, int, 1, CmpOp, VSf, IDf>(cmp_ops, q);
   passed &= test<char, float, 7, CmpOp, VSf, IDf>(cmp_ops, q);
-  passed &= test<short, double, 7, CmpOp, VSf, IDf>(cmp_ops, q);
+#ifdef ENABLE_FP64
+  passed &= test<short, double, 7, BinOp, VEf, IDf>(arith_ops, q, 1e-15);
+#endif
   passed &= test<float, float, 32, CmpOp, VSf, IDf>(cmp_ops, q);
   passed &= test<half, char, 1, CmpOp, VSf, IDf>(cmp_ops, q, 1);
   passed &= test<half, unsigned int, 32, CmpOp, VSf, IDf>(cmp_ops, q, 1);
-  passed &= test<double, half, 7, CmpOp, VSf, IDf>(cmp_ops, q);
   passed &= test<short, uint64_t, 7, CmpOp, VSf, IDf>(cmp_ops, q);
+#ifdef ENABLE_FP64
+  passed &= test<double, half, 7, CmpOp, VSf, IDf>(cmp_ops, q);
+  passed &= test<short, double, 7, CmpOp, VSf, IDf>(cmp_ops, q);
+#endif
 #ifdef USE_BF16
   passed &= test<bfloat16, int, 32, CmpOp, VSf, IDf>(cmp_ops, q);
   passed &= test<half, bfloat16, 7, CmpOp, VSf, IDf>(cmp_ops, q);

SYCL/ESIMD/api/bin_and_cmp_ops_heavy_aspect_fp64.cpp

@@ -0,0 +1,31 @@
+//==-- bin_un_cmp_ops_heavy_aspect-fp64.cpp  - DPC++ ESIMD on-device test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// Enable FP64 part of <bin_and_cmp_ops_heavy.cpp>. To be removed once DPC++
+// supports optional device features and the code could be enabled
+// unconditionally without causing failures in speculative compilation
+// of the kernels.
+//
+// REQUIRES: aspect-fp64, gpu
+// UNSUPPORTED: cuda || hip
+// TODO: esimd_emulator fails due to unimplemented 'half' type
+// XFAIL: esimd_emulator
+// RUN: %clangxx -fsycl -DENABLE_FP64 %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// Tests various binary operations applied to simd objects.
+
+// TODO
+// Arithmetic operations behaviour depends on Gen's control regiter's rounding
+// mode, which is RTNE by default:
+//    cr0.5:4 is 00b = Round to Nearest or Even (RTNE)
+// For half this leads to divergence between Gen and host (emulated) results
+// larger than certain threshold. Might need to tune the cr0 once this feature
+// is available in ESIMD.
+//
+
+#include "bin_and_cmp_ops_heavy.cpp"

SYCL/ESIMD/api/functional/ctors/ctor_converting_fp_extra.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu, level_zero
+// REQUIRES: gpu, level_zero, aspect-fp64
 // XREQUIRES: gpu
 // TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
 // The current "REQUIRES" should be replaced with "gpu" only as mentioned in

SYCL/ESIMD/api/functional/ctors/ctor_load_acc_fp_extra.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu, level_zero
+// REQUIRES: gpu, level_zero, aspect-fp64
 // XREQUIRES: gpu
 // TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
 // The current "REQUIRES" should be replaced with "gpu" only as mentioned in

SYCL/ESIMD/api/functional/ctors/ctor_load_usm_fp_extra.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu, level_zero
+// REQUIRES: gpu, level_zero, aspect-fp64
 // XREQUIRES: gpu
 // TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
 // The current "REQUIRES" should be replaced with "gpu" only as mentioned in

SYCL/ESIMD/api/functional/functions/functions_select_2d.hpp

@@ -224,8 +224,10 @@ bool run_test_for_types(sycl::queue &queue) {
   const auto great_size = get_dimensions<desired_simd_large_size>();
 #ifdef SIMD_RUN_TEST_WITH_SYCL_HALF_TYPE
   const auto all_types = get_tested_types<TestedTypes>();
-#else
+#elif SIMD_RUN_TEST_WITH_SYCL_DOUBLE_TYPE
   const auto all_types = named_type_pack<double>::generate("double");
+#else
+  const auto all_types = named_type_pack<float>::generate("float");
 #endif
 
   // Verify correctness for different select sizes.

SYCL/ESIMD/api/functional/functions/functions_select_2d_core_aspect_fp64.cpp

@@ -0,0 +1,41 @@
+//==- functions_select_2d_core_aspect_fp64.cpp - DPC++ ESIMD on-device test ==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu, level_zero, aspect-fp64
+// XREQUIRES: gpu
+// TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
+// The current "REQUIRES" should be replaced with "gpu" only as mentioned in
+// "XREQUIRES".
+// UNSUPPORTED: cuda, hip
+// RUN: %clangxx -fsycl -DSIMD_RUN_TEST_WITH_SYCL_DOUBLE_TYPE %s -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// Test for simd select for 2d function.
+// The test creates source simd instance with reference data and invokes logical
+// not operator, using core data types.
+// The test verifies that selected values can be changed with avoid to change
+// values, that hasn't beed selected.
+
+#include "functions_select_2d.hpp"
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+int main(int, char **) {
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  if (!queue.get_device().has(sycl::aspect::fp64) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  bool passed = functions::run_test_for_types<tested_types::core>(queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}