From 39cb381b88893f095a1463c026994e3aeda294fb Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Tue, 2 Dec 2025 10:44:04 -0800 Subject: [PATCH] [SYCL][ABI-breaking] Remove deprecated variadic printf implementation --- sycl/include/sycl/__spirv/spirv_ops.hpp | 14 ------- sycl/test-e2e/Basic/built-ins.cpp | 12 ------ sycl/test-e2e/DeviceLib/built-ins/printf.cpp | 16 -------- sycl/test-e2e/ESIMD/printf.cpp | 16 -------- ...nc_alloc_device_memory_reuse_zero_init.cpp | 2 - sycl/test-e2e/Printf/float.cpp | 15 ------- sycl/test/extensions/experimental-printf.cpp | 40 ------------------- sycl/test/warnings/variadic_ocl_printf.cpp | 5 --- 8 files changed, 120 deletions(-) delete mode 100644 sycl/test/extensions/experimental-printf.cpp delete mode 100644 sycl/test/warnings/variadic_ocl_printf.cpp diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp index ad0c7a31d3519..4f68506aaca03 100644 --- a/sycl/include/sycl/__spirv/spirv_ops.hpp +++ b/sycl/include/sycl/__spirv/spirv_ops.hpp @@ -838,19 +838,6 @@ __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 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, @@ -858,7 +845,6 @@ __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, template extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, Args... args); -#endif // Native builtin extension diff --git a/sycl/test-e2e/Basic/built-ins.cpp b/sycl/test-e2e/Basic/built-ins.cpp index 5ec26787c676a..2984ba302a9d4 100644 --- a/sycl/test-e2e/Basic/built-ins.cpp +++ b/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([=]() { diff --git a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp index a0e7bff0d939e..e4608f7fd483e 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp +++ b/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 #include @@ -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([=]() { @@ -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.)}} diff --git a/sycl/test-e2e/ESIMD/printf.cpp b/sycl/test-e2e/ESIMD/printf.cpp index 8f828b5e69f81..b15abe6128d26 100644 --- a/sycl/test-e2e/ESIMD/printf.cpp +++ b/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([=]() { @@ -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.)}} diff --git a/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp index e4ec248d16d43..d5c7deabc79e8 100644 --- a/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp +++ b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp @@ -9,8 +9,6 @@ #include #include -#define __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ - using T = int; void add_nodes_to_graph( exp_ext::command_graph &Graph, diff --git a/sycl/test-e2e/Printf/float.cpp b/sycl/test-e2e/Printf/float.cpp index 624a5977bc707..1fae377529cc4 100644 --- a/sycl/test-e2e/Printf/float.cpp +++ b/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([]() { do_float_test(); }); }); diff --git a/sycl/test/extensions/experimental-printf.cpp b/sycl/test/extensions/experimental-printf.cpp deleted file mode 100644 index 3efc00bbca2c9..0000000000000 --- a/sycl/test/extensions/experimental-printf.cpp +++ /dev/null @@ -1,40 +0,0 @@ -// This test is intended to check that internal -// __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-FLOAT < %t.spt -// -// 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-DOUBLE < %t.spt - -// CHECK-FLOAT: TypeFloat [[#TYPE:]] 32 -// CHECK-DOUBLE: TypeFloat [[#TYPE:]] 64 -// CHECK: Constant [[#TYPE]] [[#CONST:]] -// CHECK: ExtInst [[#]] [[#]] [[#]] printf [[#]] [[#CONST]] - -#include - -#ifdef __SYCL_DEVICE_ONLY__ -#define __SYCL_CONSTANT_AS __attribute__((opencl_constant)) -#else -#define __SYCL_CONSTANT_AS -#endif - -const __SYCL_CONSTANT_AS char fmt[] = "Hello, World! %f\n"; - -int main() { - sycl::queue q; - - q.submit([&](sycl::handler &cgh) { - cgh.single_task([=]() { - float f = 3.14; - sycl::ext::oneapi::experimental::printf(fmt, f); - }); - }); - - return 0; -} diff --git a/sycl/test/warnings/variadic_ocl_printf.cpp b/sycl/test/warnings/variadic_ocl_printf.cpp deleted file mode 100644 index 74caff358136d..0000000000000 --- a/sycl/test/warnings/variadic_ocl_printf.cpp +++ /dev/null @@ -1,5 +0,0 @@ -// RUN: %clangxx -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s - -// expected-warning@*:* {{__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ is deprecated and will be removed in a future release.}} -#include -