[SYCL] Flip default printf implementation (#13055)

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 <steffen.larsen@intel.com>
intel · Mar 20, 2024 · d704d3c · d704d3c
1 parent 43f0963
commit d704d3c
Show file tree

Hide file tree

Showing 6 changed files with 37 additions and 39 deletions.
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