[SYCL] Introduce non-variadic version of __spirv_ocl_printf (#4779)

Moved declaration of the built-in from compiler to header files, in
order to add a switch (through macro) between two versions of the
built-in: variadic and non-variadic one.

The problem with variadic version is that it follows C++ argument
promotion rules and replaces all floats with doubles, which doesn't look
good on HW, which doesn't support doubles.

Due to the fact, that we don't know the target HW in advance, we can't
selectively disable argument promotion for only some targets and doing
so will look like a hack anyway.

Therefore, a non-variadic version is introduced, which allows to avoid
argument promotion without any further changes to the compiler (and
thus, not affecting other variadic functions).

Encountering such `printf` call without argument promotion performed
should not be a surprise for a backend which consumes SPIR-V, because
SPIR-V instruction for `printf` is not variadic, all argument types are
statically known and it is legal to have both floats and doubles as
arguments at the same time. However, there is no certainty yet in proper
support of that instruction in backends: they were mostly designed for
OpenCL, where we either have argument promotion or don't have it, but
not encountering a mixed case.

Because of that, the new version is put under a macro for now, which
allows to experiment with it without disrupting existing applications
using `printf`.

Author: AlexeySachkov

clang/lib/Sema/SPIRVBuiltins.td

@@ -722,12 +722,6 @@ foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in {
 
 // 2.8. Misc instructions
 
-let IsVariadic = 1 in {
-  foreach name = ["printf"] in {
-    def : OCLSPVBuiltin<name, [Int, PointerType<ConstType<TrueChar>, ConstantAS>]>;
-  }
-}
-
 foreach name = ["prefetch"] in {
   def : OCLSPVBuiltin<name, [Void, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size]>;
 }

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -636,6 +636,16 @@ extern SYCL_EXTERNAL float __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
 __spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept;
 
+#ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+template <typename... Args>
+extern SYCL_EXTERNAL int
+__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
+                   Args... args);
+#else
+extern SYCL_EXTERNAL int
+__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
+#endif
+
 #else // if !__SYCL_DEVICE_ONLY__
 
 template <typename dataT>

sycl/test/extensions/experimental-printf.cpp

@@ -0,0 +1,40 @@
+// This test is intended to check that internal
+// __SYCL_USE_NON_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: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode -D__SYCL_USE_NON_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
+
+// CHECK-FLOAT: TypeFloat [[#TYPE:]] 32
+// CHECK-DOUBLE: TypeFloat [[#TYPE:]] 64
+// CHECK: Constant [[#TYPE]] [[#CONST:]]
+// CHECK: ExtInst [[#]] [[#]] [[#]] printf [[#]] [[#CONST]]
+
+#include <CL/sycl.hpp>
+
+#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;
+}