Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Dec 3, 2025
+0 −120
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 0 additions & 14 deletions sycl/include/sycl/__spirv/spirv_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
12 changes: 0 additions & 12 deletions sycl/test-e2e/Basic/built-ins.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand All @@ -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>([=]() {
Expand Down
16 changes: 0 additions & 16 deletions sycl/test-e2e/DeviceLib/built-ins/printf.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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>([=]() {
Expand All @@ -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.)}}

Expand Down
16 changes: 0 additions & 16 deletions sycl/test-e2e/ESIMD/printf.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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>([=]() {
Expand All @@ -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.)}}

Expand Down
2 changes: 0 additions & 2 deletions sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_device_memory_reuse_zero_init.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
15 changes: 0 additions & 15 deletions sycl/test-e2e/Printf/float.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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(); });
});
Expand Down
40 changes: 0 additions & 40 deletions sycl/test/extensions/experimental-printf.cpp
View file
Open in desktop

This file was deleted.

5 changes: 0 additions & 5 deletions sycl/test/warnings/variadic_ocl_printf.cpp
View file
Open in desktop

This file was deleted.

Loading