From d704d3c979d72f2282f076ab4be91570171876c8 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 20 Mar 2024 08:18:38 +0100 Subject: [PATCH] [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 --- sycl/include/CL/__spirv/spirv_ops.hpp | 10 +++++----- sycl/test-e2e/Basic/built-ins.cpp | 13 ++++++------- sycl/test-e2e/DeviceLib/built-ins/printf.cpp | 18 +++++++++--------- sycl/test-e2e/ESIMD/printf.cpp | 18 +++++++++--------- sycl/test-e2e/Printf/float.cpp | 9 ++++----- sycl/test/extensions/experimental-printf.cpp | 8 ++++---- 6 files changed, 37 insertions(+), 39 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 5dcb2ff056921..f477fe309045d 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -1183,7 +1183,11 @@ __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 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, @@ -1191,10 +1195,6 @@ __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, template 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 index 6e8b545d1e6a6..995699e9d6ef7 100644 --- 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 index fcbb2dd029b3c..718d9b292044b 100644 --- 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 @@ -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([=]() { @@ -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 index a14fa0a418cf4..181fc2dd02824 100644 --- 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([=]() { @@ -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 index a79329f91bd03..e2c8caf42979b 100644 --- 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([]() { do_float_test(); }); }); diff --git a/sycl/test/extensions/experimental-printf.cpp b/sycl/test/extensions/experimental-printf.cpp index 4c9269e54495a..3efc00bbca2c9 100644 --- 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