[SYCL] Add support for NVPTX device printf (#4293)

Use `::printf` when not compiling for `__SPIR__`, this allows the use of
`EmitNVPTXDevicePrintfCallExpr` which packs the var args and dispatches
to CUDA's `vprintf`.

Fixes #1154

Author: jchlanda

clang/lib/CodeGen/CGGPUBuiltin.cpp

@@ -118,8 +118,19 @@ CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E,
 
   // Invoke vprintf and return.
   llvm::Function* VprintfFunc = GetVprintfDeclaration(CGM.getModule());
-  return RValue::get(Builder.CreateCall(
-      VprintfFunc, {Args[0].getRValue(*this).getScalarVal(), BufferPtr}));
+  auto FormatSpecifier = Args[0].getRValue(*this).getScalarVal();
+  // Check if the format specifier is in the constant address space, vprintf is
+  // oblivious to address spaces, so it would have to be casted away.
+  if (Args[0]
+          .getRValue(*this)
+          .getScalarVal()
+          ->getType()
+          ->getPointerAddressSpace() == 4)
+    FormatSpecifier = Builder.CreateAddrSpaceCast(
+        FormatSpecifier, llvm::Type::getInt8PtrTy(Ctx));
+
+  return RValue::get(
+      Builder.CreateCall(VprintfFunc, {FormatSpecifier, BufferPtr}));
 }
 
 RValue

clang/lib/Sema/SemaSYCL.cpp

@@ -420,6 +420,8 @@ static bool IsSyclMathFunc(unsigned BuiltinID) {
 bool Sema::isKnownGoodSYCLDecl(const Decl *D) {
   if (const FunctionDecl *FD = dyn_cast<FunctionDecl>(D)) {
     const IdentifierInfo *II = FD->getIdentifier();
+    if (FD->getBuiltinID() == Builtin::BIprintf)
+      return true;
     const DeclContext *DC = FD->getDeclContext();
     if (II && II->isStr("__spirv_ocl_printf") &&
         !FD->isDefined() &&

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,12 @@
+// 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"
+
+static const __SYCL_CONSTANT_AS 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); });
+}

sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

@@ -61,11 +61,11 @@ namespace experimental {
 //
 template <typename... Args>
 int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) {
-#ifdef __SYCL_DEVICE_ONLY__
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
   return __spirv_ocl_printf(__format, args...);
 #else
   return ::printf(__format, args...);
-#endif
+#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
 }
 
 } // namespace experimental