[SYCL] Add NVPTX printf test and printf support in mock SYCL header

Author: jchlanda

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -2,6 +2,8 @@
 
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 
+extern "C" int printf(const char* fmt, ...);
+
 // Dummy runtime classes to model SYCL API.
 inline namespace cl {
 namespace sycl {
@@ -310,6 +312,21 @@ class spec_constant {
     return get();
   }
 };
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define __SYCL_CONSTANT_AS __attribute__((opencl_constant))
+#else
+#define __SYCL_CONSTANT_AS
+#endif
+template <typename... Args>
+int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) {
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
+  return __spirv_ocl_printf(__format, args...);
+#else
+  return ::printf(__format, args...);
+#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
+}
+
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext

clang/test/CodeGenSYCL/nvptx-printf.cpp

@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-nvidia-cuda-sycldevice -std=c++11 -S -emit-llvm -x c++ %s -o - | FileCheck %s
+
+#include "Inputs/sycl.hpp"
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define CONSTANT __attribute__((opencl_constant))
+#else
+#define CONSTANT
+#endif
+
+static const CONSTANT char format_2[] = "Hello! %d %f\n";
+
+int main() {
+  // Make sure that device printf is dispatched to CUDA's vprintf syscall.
+  // CHECK: alloca %printf_args
+  // CHECK: call i32 @vprintf
+  cl::sycl::kernel_single_task<class first_kernel>([]() { cl::sycl::ext::oneapi::experimental::printf(format_2, 123, 1.23); });
+}