[SYCL] Add 'cl::sycl::intel::experimental::printf'

sycl/include/CL/sycl.hpp

@@ -21,6 +21,7 @@
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/image.hpp>
+#include <CL/sycl/intel/builtins.hpp>
 #include <CL/sycl/intel/function_pointer.hpp>
 #include <CL/sycl/intel/sub_group.hpp>
 #include <CL/sycl/item.hpp>

sycl/include/CL/sycl/intel/builtins.hpp

@@ -0,0 +1,77 @@
+//==------ builtins.hpp - Non-standard SYCL built-in functions -------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define CONSTANT_AS __attribute__((opencl_constant))
+// Note: __format string is declared in constant address space to be compatible
+// with OpenCL C
+extern int __spirv_ocl_printf(const CONSTANT_AS char *__format, ...);
+#else
+#define CONSTANT_AS
+#endif
+
+namespace cl {
+namespace sycl {
+namespace intel {
+namespace experimental {
+
+// Provides functionality to print data from kernels in a C way:
+// - On non-host devices this function is directly mapped to printf from
+//   OpenCL C
+// - On host device, this function should be equivalent to standard printf
+//   function from C/C++.
+//
+// Please refer to corresponding section in OpenCL C specification to find
+// information about format string and its differences from standard C rules.
+//
+// This function is placed under 'experimental' namespace on purpose, because it
+// has too much caveats you need to be aware of before using it. Please find
+// them below and read carefully before using it:
+//
+// - According to the OpenCL spec, the format string must be
+// resolvable at compile time i.e. cannot be dynamically created by the
+// executing program.
+//
+// - According to the OpenCL spec, the format string must reside in constant
+// address space. This requires to perform "tricky" declarations of them, see
+// test/built-ins/printf.cpp for examples
+// FIXME: this potentially can be done on SYCL FE side automatically
+//
+// - The format string is interpreted according to the OpenCL C spec, where all
+// data types has fixed size, opposed to C++ types which doesn't guarantee
+// the exact width of particular data types (except, may be, char). This might
+// lead to unexpected result, for example: %ld in OpenCL C means that printed
+// argument has 'long' type which is 64-bit wide by the OpenCL C spec. However,
+// by C++ spec long is just at least 32-bit wide, so, you need to ensure (by
+// performing a cast, for example) that if you use %ld specifier, you pass
+// 64-bit argument to the cl::sycl::experimental::printf
+//
+// - OpenCL spec defines several additional features, like, for example, 'v'
+// modifier which allows to print OpenCL vectors: note that these features are
+// not available on host device and therefore their usage should be either
+// guarded using __SYCL_DEVICE_ONLY__ preprocessor macro or avoided in favor
+// of more portable solutions if needed
+//
+template <typename... Args>
+int printf(const CONSTANT_AS char *__format, Args... args) {
+#ifdef __SYCL_DEVICE_ONLY__
+  return __spirv_ocl_printf(__format, args...);
+#else
+  return ::printf(__format, args...);
+#endif
+}
+
+} // namespace experimental
+} // namespace intel
+} // namespace sycl
+} // namespace cl
+
+#undef CONSTANT_AS

sycl/test/built-ins/printf.cpp

@@ -0,0 +1,131 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out | FileCheck %s
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+
+#include <CL/sycl.hpp>
+
+#include <cstdint>
+#include <iostream>
+
+using namespace cl::sycl;
+
+// According to OpenCL C spec, the format string must be in constant address
+// space
+#ifdef __SYCL_DEVICE_ONLY__
+#define CONSTANT __attribute__((opencl_constant))
+#else
+#define CONSTANT
+#endif
+
+// This is one of the possible ways to define a format string in a correct
+// address space
+static const CONSTANT char format_hello_world[] = "Hello, World!\n";
+
+// Static isn't really needed if you define it in global scope
+const CONSTANT char format_int[] = "%d\n";
+
+static const CONSTANT char format_vec[] = "%d,%d,%d,%d\n";
+
+const CONSTANT char format_hello_world_2[] = "%lu: Hello, World!\n";
+
+int main() {
+  {
+    default_selector Selector;
+    queue Queue(Selector);
+
+    Queue.submit([&](handler &CGH) {
+      CGH.single_task<class integral>([=]() {
+        // String
+        intel::experimental::printf(format_hello_world);
+        // Due to a bug in Intel CPU Runtime for OpenCL on Windows, information
+        // printed using such format strings (without %-specifiers) might
+        // appear in different order if output is redirected to a file or
+        // another app
+        // FIXME: strictly check output order once the bug is fixed
+        // CHECK: {{(Hello, World!)?}}
+
+        // Integral types
+        intel::experimental::printf(format_int, (int32_t)123);
+        intel::experimental::printf(format_int, (int32_t)-123);
+        // CHECK: 123
+        // CHECK-NEXT: -123
+
+        // Floating point types
+        {
+          // You can declare format string in non-global scope, but in this case
+          // static keyword is required
+          static const CONSTANT char format[] = "%f\n";
+          intel::experimental::printf(format, 33.4f);
+          intel::experimental::printf(format, -33.4f);
+        }
+        // CHECK-NEXT: 33.4
+        // CHECK-NEXT: -33.4
+
+        // Vectors
+        cl::sycl::vec<int, 4> v4{5, 6, 7, 8};
+#ifdef __SYCL_DEVICE_ONLY__
+        // On device side, vectors can be printed via native OpenCL types:
+        using ocl_int4 = cl::sycl::vec<int, 4>::vector_t;
+        {
+          static const CONSTANT char format[] = "%v4d\n";
+          intel::experimental::printf(format, (ocl_int4)v4);
+        }
+
+        // However, you are still able to print them by-element:
+        {
+          intel::experimental::printf(format_vec, (int32_t)v4.w(),
+                                      (int32_t)v4.z(), (int32_t)v4.y(),
+                                      (int32_t)v4.x());
+        }
+#else
+        // On host side you always have to print them by-element:
+        intel::experimental::printf(format_vec, (int32_t)v4.x(),
+                                    (int32_t)v4.y(), (int32_t)v4.z(),
+                                    (int32_t)v4.w());
+        intel::experimental::printf(format_vec, (int32_t)v4.w(),
+                                    (int32_t)v4.z(), (int32_t)v4.y(),
+                                    (int32_t)v4.x());
+#endif // __SYCL_DEVICE_ONLY__
+       // CHECK-NEXT: 5,6,7,8
+       // CHECK-NEXT: 8,7,6,5
+
+        // Pointers
+        int a = 5;
+        int *Ptr = &a;
+        // According to OpenCL spec, argument should be a void pointer
+        {
+          static const CONSTANT char format[] = "%p\n";
+          intel::experimental::printf(format, (void *)Ptr);
+        }
+        // CHECK-NEXT: {{(0x)?[0-9a-fA-F]+$}}
+      });
+    });
+    Queue.wait();
+
+    // printf in parallel_for
+    Queue.submit([&](handler &CGH) {
+      CGH.parallel_for<class stream_string>(range<1>(10), [=](id<1> i) {
+        // cast to uint64_t to be sure that we pass 64-bit unsigned value
+        intel::experimental::printf(format_hello_world_2, (uint64_t)i.get(0));
+      });
+    });
+    Queue.wait();
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+    // CHECK-NEXT: {{[0-9]+}}: Hello, World!
+  }
+
+// FIXME: strictly check output order once the bug mentioned above is fixed
+// CHECK: {{(Hello, World!)?}}
+
+  return 0;
+}