[SYCL] Flip default printf implementation (#13055)

steffenlarsen · web-flow · commit d704d3c979d7 · 2024-03-20T08:18:38.000+01:00
Currently, the implementation of printf uses a variadic implementation
by default. This proved to be problematic as the implementation would
promote float arguments to doubles, implicitly requring fp64. As a
result of this, an alternative implementation was introduced, but was
made enableable using a __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
directive to avoid problems with targets that did not support it. We
expect most relevant backends to support it now, so we flip the default.

---------

Signed-off-by: Larsen, Steffen &lt;steffen.larsen@intel.com&gt;
diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -1183,18 +1183,18 @@ __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
 
-#ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
+extern __DPCPP_SYCL_EXTERNAL int
+__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
+extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
+#else
 template <typename... Args>
 extern __DPCPP_SYCL_EXTERNAL int
 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
                    Args... args);
 template <typename... Args>
 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format,
                                                     Args... args);
-#else
-extern __DPCPP_SYCL_EXTERNAL int
-__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
-extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
 #endif
 
 // Native builtin extension
diff --git a/sycl/test-e2e/Basic/built-ins.cpp b/sycl/test-e2e/Basic/built-ins.cpp
@@ -1,8 +1,8 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
 
-// RUN: %{build} -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_nonvar.out
-// RUN: %{run} %t_nonvar.out | FileCheck %s
+// RUN: %{build} -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_var.out
+// RUN: %{run} %t_var.out | FileCheck %s
 
 // Hits an assertion with AMD:
 // XFAIL: hip_amd
@@ -26,12 +26,11 @@ static const CONSTANT char format[] = "Hello, World! %d %f\n";
 int main() {
   s::queue q{};
 
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_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;
+    std::cout << "Test with __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ defined is "
+                 "skipped because the device did not have fp64."
+              << std::endl;
     return 0;
   }
 #endif
diff --git a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp
@@ -2,11 +2,11 @@
 // HIP doesn't support printf.
 // CUDA doesn't support vector format specifiers ("%v").
 //
-// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
 //
-// RUN: %{build} -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_nonvar.out
-// RUN: %{run} %t_nonvar.out | FileCheck %s
+// RUN: %{build} -fsycl-device-code-split=per_kernel -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_var.out
+// RUN: %{run} %t_var.out | FileCheck %s
 
 #include <sycl/sycl.hpp>
 
@@ -96,13 +96,13 @@ int main() {
     Queue.wait();
   }
 
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   // Currently printf will promote floating point values to doubles.
-  // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ changes the behavior to not use
-  // a variadic function, so if it is defined it will not promote the floating
+  // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ changes the behavior to use
+  // a variadic function, so if it is defined it will promote the floating
   // point arguments.
   if (Queue.get_device().has(sycl::aspect::fp64))
-#endif // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   {
     Queue.submit([&](handler &CGH) {
       CGH.single_task<class floating_points>([=]() {
@@ -118,12 +118,12 @@ int main() {
     });
     Queue.wait();
   }
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   else {
     std::cout << "Skipped floating point test." << std::endl;
     std::cout << "Skipped floating point test." << std::endl;
   }
-#endif // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   // CHECK-NEXT: {{(33.4|Skipped floating point test.)}}
   // CHECK-NEXT: {{(-33.4|Skipped floating point test.)}}
 
diff --git a/sycl/test-e2e/ESIMD/printf.cpp b/sycl/test-e2e/ESIMD/printf.cpp
@@ -7,11 +7,11 @@
 //===----------------------------------------------------------------------===//
 //
 //
-// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
 //
-// RUN: %{build} -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_nonvar.out
-// RUN: %{run} %t_nonvar.out | FileCheck %s
+// RUN: %{build} -fsycl-device-code-split=per_kernel -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -o %t_var.out
+// RUN: %{run} %t_var.out | FileCheck %s
 //
 //===----------------------------------------------------------------------===//
 //
@@ -70,13 +70,13 @@ int main() {
     Queue.wait();
   }
 
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   // Currently printf will promote floating point values to doubles.
-  // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ changes the behavior to not use
-  // a variadic function, so if it is defined it will not promote the floating
+  // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ changes the behavior to use
+  // a variadic function, so if it is defined it will promote the floating
   // point arguments.
   if (Queue.get_device().has(sycl::aspect::fp64))
-#endif // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   {
     Queue.submit([&](handler &CGH) {
       CGH.single_task<class floating_points>([=]() {
@@ -92,12 +92,12 @@ int main() {
     });
     Queue.wait();
   }
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   else {
     std::cout << "Skipped floating point test." << std::endl;
     std::cout << "Skipped floating point test." << std::endl;
   }
-#endif // __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   // CHECK-NEXT: {{(33.4|Skipped floating point test.)}}
   // CHECK-NEXT: {{(-33.4|Skipped floating point test.)}}
 
diff --git a/sycl/test-e2e/Printf/float.cpp b/sycl/test-e2e/Printf/float.cpp
@@ -9,9 +9,8 @@
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
-// FIXME: Remove dedicated non-variadic printf testing once the headers
-//        enforce it by default.
-// RUN: %{build} -o %t.nonvar.out -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+// FIXME: Remove dedicated variadic printf testing once the option is removed.
+// RUN: %{build} -o %t.nonvar.out -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
 // RUN: %{run} %t.nonvar.out | FileCheck %s
 // FIXME: Remove dedicated constant address space testing once generic AS
 //        support is considered stable.
@@ -48,7 +47,7 @@ class FloatTest;
 int main() {
   queue q;
 
-#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   if (!q.get_device().has(aspect::fp64)) {
     std::cout << "Skipping the actual test due to variadic argument promotion. "
                  "Printing hard-coded output from the host side:\n"
@@ -58,7 +57,7 @@ int main() {
               << std::endl;
     return 0;
   }
-#endif // !__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   q.submit([](handler &cgh) {
     cgh.single_task<FloatTest>([]() { do_float_test(); });
   });
diff --git a/sycl/test/extensions/experimental-printf.cpp b/sycl/test/extensions/experimental-printf.cpp
@@ -1,15 +1,15 @@
 // This test is intended to check that internal
-// __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ works as expected, i.e. we can
+// __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ works as expected, i.e. we can
 // see printf ExtInst regardless of the macro presence and that argument
 // promotion is disabled if the macro is present.
 //
 // RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode %s -o %t.spv
 // RUN: llvm-spirv -to-text %t.spv -o %t.spt
-// RUN: FileCheck %s --check-prefixes CHECK,CHECK-DOUBLE < %t.spt
+// RUN: FileCheck %s --check-prefixes CHECK,CHECK-FLOAT < %t.spt
 //
-// RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ %s -o %t.spv
+// RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ %s -o %t.spv
 // RUN: llvm-spirv -to-text %t.spv -o %t.spt
-// RUN: FileCheck %s --check-prefixes CHECK,CHECK-FLOAT < %t.spt
+// RUN: FileCheck %s --check-prefixes CHECK,CHECK-DOUBLE < %t.spt
 
 // CHECK-FLOAT: TypeFloat [[#TYPE:]] 32
 // CHECK-DOUBLE: TypeFloat [[#TYPE:]] 64
