[SYCL] Prevent use of fp64 and fp16 when unsupported in more tests

SYCL/Basic/bit_cast/bit_cast.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl -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,11 +13,11 @@ constexpr sycl::access::mode sycl_write = sycl::access::mode::write;
 
 template <typename To, typename From> class BitCastKernel;
 
-template <typename To, typename From> To doBitCast(const From &ValueToConvert) {
+template <typename To, typename From>
+To doBitCast(sycl::queue Queue, const From &ValueToConvert) {
   std::vector<To> Vec(1);
   {
     sycl::buffer<To, 1> Buf(Vec.data(), 1);
-    sycl::queue Queue;
     Queue.submit([&](sycl::handler &cgh) {
       auto acc = Buf.template get_access<sycl_write>(cgh);
       cgh.single_task<class BitCastKernel<To, From>>([=]() {
@@ -28,8 +28,10 @@ template <typename To, typename From> To doBitCast(const From &ValueToConvert) {
   return Vec[0];
 }
 
-template <typename To, typename From> int test(const From &Value) {
-  auto ValueConvertedTwoTimes = doBitCast<From>(doBitCast<To>(Value));
+template <typename To, typename From>
+int test(sycl::queue Queue, const From &Value) {
+  auto ValueConvertedTwoTimes =
+      doBitCast<From>(Queue, doBitCast<To>(Queue, Value));
   bool isOriginalValueEqualsToConvertedTwoTimes = false;
   if (std::is_integral<From>::value) {
     isOriginalValueEqualsToConvertedTwoTimes = Value == ValueConvertedTwoTimes;
@@ -54,31 +56,34 @@ template <typename To, typename From> int test(const From &Value) {
 }
 
 int main() {
+  sycl::queue Queue;
   int ReturnCode = 0;
 
-  std::cout << "sycl::half to unsigned short ...\n";
-  ReturnCode += test<unsigned short>(sycl::half(1.0f));
+  if (Queue.get_device().has(sycl::aspect::fp16)) {
+    std::cout << "sycl::half to unsigned short ...\n";
+    ReturnCode += test<unsigned short>(Queue, sycl::half(1.0f));
 
-  std::cout << "unsigned short to sycl::half ...\n";
-  ReturnCode += test<sycl::half>(static_cast<unsigned short>(16384));
+    std::cout << "unsigned short to sycl::half ...\n";
+    ReturnCode += test<sycl::half>(Queue, static_cast<unsigned short>(16384));
 
-  std::cout << "sycl::half to short ...\n";
-  ReturnCode += test<short>(sycl::half(1.0f));
+    std::cout << "sycl::half to short ...\n";
+    ReturnCode += test<short>(Queue, sycl::half(1.0f));
 
-  std::cout << "short to sycl::half ...\n";
-  ReturnCode += test<sycl::half>(static_cast<short>(16384));
+    std::cout << "short to sycl::half ...\n";
+    ReturnCode += test<sycl::half>(Queue, static_cast<short>(16384));
+  }
 
   std::cout << "int to float ...\n";
-  ReturnCode += test<float>(static_cast<int>(2));
+  ReturnCode += test<float>(Queue, static_cast<int>(2));
 
   std::cout << "float to int ...\n";
-  ReturnCode += test<int>(static_cast<float>(-2.4f));
+  ReturnCode += test<int>(Queue, static_cast<float>(-2.4f));
 
   std::cout << "unsigned int to float ...\n";
-  ReturnCode += test<float>(static_cast<unsigned int>(6));
+  ReturnCode += test<float>(Queue, static_cast<unsigned int>(6));
 
   std::cout << "float to unsigned int ...\n";
-  ReturnCode += test<unsigned int>(static_cast<float>(-2.4f));
+  ReturnCode += test<unsigned int>(Queue, static_cast<float>(-2.4f));
 
   return ReturnCode;
 }

SYCL/Basic/built-ins.cpp

@@ -3,6 +3,11 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ %s -o %t_nonvar.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_nonvar.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t_nonvar.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t_nonvar.out %ACC_CHECK_PLACEHOLDER
+
 // CUDA does not support printf.
 // UNSUPPORTED: cuda
 //
@@ -28,15 +33,25 @@ static const CONSTANT char format[] = "Hello, World! %d %f\n";
 int main() {
   s::queue q{};
 
+#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+  if (!q.get_device().has(sycl::aspect::fp64)) {
+    std::cout
+        << "Test without __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ defined is "
+           "skipped because the device did not have fp64."
+        << std::endl;
+    return 0;
+  }
+#endif
+
   // Test printf
   q.submit([&](s::handler &CGH) {
      CGH.single_task<class printf>([=]() {
-       s::ext::oneapi::experimental::printf(format, 123, 1.23);
+       s::ext::oneapi::experimental::printf(format, 123, 1.23f);
        // CHECK: {{(Hello, World! 123 1.23)?}}
      });
    }).wait();
 
-  s::ext::oneapi::experimental::printf(format, 321, 3.21);
+  s::ext::oneapi::experimental::printf(format, 321, 3.21f);
   // CHECK: {{(Hello, World! 123 1.23)?}}
 
   // Test common
@@ -47,7 +62,7 @@ int main() {
       auto AccMin = BufMin.get_access<s::access::mode::write>(cgh);
       auto AccMax = BufMax.get_access<s::access::mode::write>(cgh);
       cgh.single_task<class common>([=]() {
-        AccMax[0] = s::max(s::cl_float2{0.5f, 2.5}, s::cl_float2{2.3f, 2.3});
+        AccMax[0] = s::max(s::cl_float2{0.5f, 2.5f}, s::cl_float2{2.3f, 2.3f});
         AccMin[0] = s::min(s::cl_float{0.5f}, s::cl_float{2.3f});
       });
     });

SYCL/Basic/half_builtins.cpp

@@ -165,6 +165,14 @@ template <int N> bool check(vec<float, N> a, vec<float, N> b) {
 
 int main() {
   queue q;
+
+  if (!q.get_device().has(sycl::aspect::fp16)) {
+    std::cout
+        << "Test was skipped because the selected device does not support fp16"
+        << std::endl;
+    return 0;
+  }
+
   float16 a, b, c, d;
   for (int i = 0; i < SZ_max; i++) {
     a[i] = i / (float)SZ_max;

SYCL/Basic/scalar_vec_access.cpp

@@ -29,7 +29,7 @@ int main() {
       // Test that it is possible to get a reference to single element of the
       // vector type. This behavior could possibly change in the future, this
       // test is necessary to track that.
-      float4_t my_float4 = {0.0, 1.0, 2.0, 3.0};
+      float4_t my_float4 = {0.0f, 1.0f, 2.0f, 3.0f};
       float f[4];
       for (int i = 0; i < 4; ++i) {
         f[i] = reinterpret_cast<float *>(&my_float4)[i];
@@ -40,14 +40,14 @@ int main() {
       }
 
       // Test that there is no template resolution error
-      sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
+      sycl::float4 a = {1.0f, 2.0f, 3.0f, 4.0f};
       out << sycl::native::recip(a.x()) << sycl::endl;
     });
   });
   Q.wait();
 
   // Test that there is no ambiguity in overload resolution.
-  sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
+  sycl::float4 a = {1.0f, 2.0f, 3.0f, 4.0f};
   std::cout << a.x() << std::endl;
 
   return 0;

SYCL/Basic/stream/stream.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 %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
@@ -97,26 +97,14 @@ int main() {
         // CHECK-NEXT: -12345678901245
         // CHECK-NEXT: 12345678901245
 
-        // Floating point types
+        // Floats
         Out << 33.4f << endl;
-        Out << 5.2 << endl;
         Out << -33.4f << endl;
-        Out << -5.2 << endl;
-        Out << 0.0003 << endl;
-        Out << -1.0 / 0.0 << endl;
-        Out << 1.0 / 0.0 << endl;
-        Out << sycl::sqrt(-1.0) << endl;
         Out << -1.0f / 0.0f << endl;
         Out << 1.0f / 0.0f << endl;
         Out << sycl::sqrt(-1.0f) << endl;
         // CHECK-NEXT: 33.4
-        // CHECK-NEXT: 5.2
         // CHECK-NEXT: -33.4
-        // CHECK-NEXT: -5.2
-        // CHECK-NEXT: 0.0003
-        // CHECK-NEXT: -inf
-        // CHECK-NEXT: inf
-        // CHECK-NEXT: nan
         // CHECK-NEXT: -inf
         // CHECK-NEXT: inf
         // CHECK-NEXT: nan
@@ -205,6 +193,33 @@ int main() {
     });
     Queue.wait();
 
+    if (Queue.get_device().has(sycl::aspect::fp64)) {
+      Queue.submit([&](handler &CGH) {
+        stream Out(1024, 80, CGH);
+        CGH.single_task<class doubles>([=]() {
+          // Double
+          Out << 5.2 << endl;
+          Out << -5.2 << endl;
+          Out << 0.0003 << endl;
+          Out << -1.0 / 0.0 << endl;
+          Out << 1.0 / 0.0 << endl;
+          Out << sycl::sqrt(-1.0) << endl;
+        });
+      });
+      Queue.wait();
+    } else {
+      // Repeat skipped message same number of times as the number of skipped
+      // output lines.
+      for (size_t I = 0; I < 6; ++I)
+        std::cout << "Skipped double test." << std::endl;
+    }
+    // CHECK-NEXT: {{(5.2|Skipped double test.)}}
+    // CHECK-NEXT: {{(-5.2|Skipped double test.)}}
+    // CHECK-NEXT: {{(0.0003|Skipped double test.)}}
+    // CHECK-NEXT: {{(-inf|Skipped double test.)}}
+    // CHECK-NEXT: {{(inf|Skipped double test.)}}
+    // CHECK-NEXT: {{(nan|Skipped double test.)}}
+
     // Stream in parallel_for
     Queue.submit([&](handler &CGH) {
       stream Out(1024, 80, CGH);

SYCL/Basic/vector_operators.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl -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
@@ -38,6 +38,7 @@ template <typename T, int N> void check_vector_size() {
 }
 
 int main() {
+  s::queue Queue;
 
   /* Separate checks for NumElements=1 edge case */
 
@@ -46,7 +47,6 @@ int main() {
     vec_type res;
     {
       s::buffer<vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isequal_vec_op_1_elem>([=]() {
@@ -67,7 +67,6 @@ int main() {
     vec_type res;
     {
       s::buffer<vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isequal_vec_op_1_elem_scalar>([=]() {
@@ -94,7 +93,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isequal_vec_op>([=]() {
@@ -109,12 +107,11 @@ int main() {
   }
 
   // Operator <, cl_double
-  {
+  if (Queue.get_device().has(sycl::aspect::fp64)) {
     using res_vec_type = s::vec<s::cl_long, 4>;
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isless_vec_op>([=]() {
@@ -134,7 +131,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isgreater_vec_op>([=]() {
@@ -149,12 +145,11 @@ int main() {
   }
 
   // Operator <=, cl_half
-  {
+  if (Queue.get_device().has(sycl::aspect::fp16)) {
     using res_vec_type = s::vec<s::cl_short, 4>;
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isnotgreater_vec_op>([=]() {
@@ -176,7 +171,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isnotless_vec_op>([=]() {
@@ -196,7 +190,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class isnotequal_vec_op>([=]() {
@@ -216,7 +209,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class logical_and_vec_op>([=]() {
@@ -236,7 +228,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class logical_or_vec_op>([=]() {
@@ -257,7 +248,6 @@ int main() {
     res_vec_type res;
     {
       s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
-      s::queue Queue;
       Queue.submit([&](s::handler &cgh) {
         auto Acc = Buf.get_access<s::access::mode::write>(cgh);
         cgh.single_task<class as_op>([=]() {