[SYCL][ABI-break] Remove deprecated variadic printf implementation (#20800)

Co-authored-by: Steffen Larsen [email protected]

Author: againull

sycl/include/sycl/__spirv/spirv_ops.hpp

@@ -838,27 +838,13 @@ __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
 
-#if defined(__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__) &&                         \
-    !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
-#if defined(__clang__)
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wpedantic"
-#warning                                                                       \
-    "__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ is deprecated and will be removed in a future release."
-#pragma clang diagnostic pop
-#endif
-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);
-#endif
 
 // Native builtin extension
 

sycl/test-e2e/Basic/built-ins.cpp

@@ -1,9 +1,6 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
 
-// RUN: %{build} -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -Wno-#warnings -o %t_var.out
-// RUN: %{run} %t_var.out | FileCheck %s
-
 // Hits an assertion and kernel page fault with AMD:
 // UNSUPPORTED: target-amd
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/14404
@@ -28,15 +25,6 @@ static const CONSTANT char format[] = "Hello, World! %d %f\n";
 int main() {
   s::queue q{};
 
-#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
-  if (!q.get_device().has(sycl::aspect::fp64)) {
-    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
-
   // Test printf
   q.submit([&](s::handler &CGH) {
      CGH.single_task<class printf>([=]() {

sycl/test-e2e/DeviceLib/built-ins/printf.cpp

@@ -4,9 +4,6 @@
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
-//
-// RUN: %{build} -fsycl-device-code-split=per_kernel -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -Wno-#warnings -o %t_var.out
-// RUN: %{run} %t_var.out | FileCheck %s
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/builtins.hpp>
@@ -98,13 +95,6 @@ int main() {
     Queue.wait();
   }
 
-#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
-  // Currently printf will promote floating point values to doubles.
-  // __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_VARIADIC_SPIRV_OCL_PRINTF__
   {
     Queue.submit([&](handler &CGH) {
       CGH.single_task<class floating_points>([=]() {
@@ -120,12 +110,6 @@ int main() {
     });
     Queue.wait();
   }
-#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_VARIADIC_SPIRV_OCL_PRINTF__
   // CHECK-NEXT: {{(33.4|Skipped floating point test.)}}
   // CHECK-NEXT: {{(-33.4|Skipped floating point test.)}}
 

sycl/test-e2e/ESIMD/printf.cpp

@@ -10,9 +10,6 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
 //
-// RUN: %{build} -fsycl-device-code-split=per_kernel -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -Wno-#warnings -o %t_var.out
-// RUN: %{run} %t_var.out | FileCheck %s
-//
 //===----------------------------------------------------------------------===//
 //
 // The test checks that ESIMD kernels support printf functionality.
@@ -67,13 +64,6 @@ int main() {
     Queue.wait();
   }
 
-#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
-  // Currently printf will promote floating point values to doubles.
-  // __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_VARIADIC_SPIRV_OCL_PRINTF__
   {
     Queue.submit([&](handler &CGH) {
       CGH.single_task<class floating_points>([=]() {
@@ -89,12 +79,6 @@ int main() {
     });
     Queue.wait();
   }
-#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_VARIADIC_SPIRV_OCL_PRINTF__
   // CHECK-NEXT: {{(33.4|Skipped floating point test.)}}
   // CHECK-NEXT: {{(-33.4|Skipped floating point test.)}}
 

sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp

@@ -9,8 +9,6 @@
 #include <sycl/ext/oneapi/experimental/async_alloc/memory_pool.hpp>
 #include <sycl/ext/oneapi/experimental/async_alloc/memory_pool_properties.hpp>
 
-#define __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
-
 using T = int;
 void add_nodes_to_graph(
     exp_ext::command_graph<exp_ext::graph_state::modifiable> &Graph,

sycl/test-e2e/Printf/float.cpp

@@ -8,9 +8,6 @@
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
-// FIXME: Remove dedicated variadic printf testing once the option is removed.
-// RUN: %{build} -o %t.nonvar.out -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -Wno-#warnings
-// RUN: %{run} %t.nonvar.out | FileCheck %s
 // FIXME: Remove dedicated constant address space testing once generic AS
 //        support is considered stable.
 // RUN: %{build} -o %t.constant.out -DTEST_CONSTANT_AS
@@ -46,18 +43,6 @@ class FloatTest;
 
 int main() {
   queue q;
-
-#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"
-              << "3.140000e+00, 3.140000E+00\n"
-                 "0x1.91eb86p+1, 0X1.91EB86P+1\n"
-                 "3.14, 3.14"
-              << std::endl;
-    return 0;
-  }
-#endif // __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
   q.submit([](handler &cgh) {
     cgh.single_task<FloatTest>([]() { do_float_test(); });
   });