[Clang][OpenCL][NVPTX] using vprintf to implement builtin printf on OpenCL with NVPTX

Author: boxu-zhang

clang/lib/AST/ASTContext.cpp

@@ -11481,6 +11481,11 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
     // FIXME: There's no way to have a built-in with an rvalue ref arg.
     case 'C':
       Type = Type.withConst();
+      // adjust 'const char *' to 'const char __constant *' on OpenCL
+      if (Context.getLangOpts().OpenCL &&
+          Type.getTypePtr() == Context.CharTy.getTypePtr()) {
+        Type = Context.getAddrSpaceQualType(Type, LangAS::opencl_constant);
+      }
       break;
     case 'D':
       Type = Context.getVolatileType(Type);

clang/lib/AST/Decl.cpp

@@ -3600,8 +3600,10 @@ unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const {
 
   // OpenCL v1.2 s6.9.f - The library functions defined in
   // the C99 standard headers are not available.
+  // EXCEPTION: printf is supported for AMDGPU
   if (Context.getLangOpts().OpenCL &&
-      Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID))
+      Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID) &&
+      BuiltinID != Builtin::BIprintf)
     return 0;
 
   // CUDA does not have device-side standard library. printf and malloc are the

clang/lib/CodeGen/CGGPUBuiltin.cpp

@@ -135,6 +135,7 @@ bool containsNonScalarVarargs(CodeGenFunction *CGF, const CallArgList &Args) {
 RValue EmitDevicePrintfCallExpr(const CallExpr *E, CodeGenFunction *CGF,
                                 llvm::Function *Decl, bool WithSizeArg) {
   CodeGenModule &CGM = CGF->CGM;
+  llvm::LLVMContext &Ctx = CGM.getLLVMContext();
   CGBuilderTy &Builder = CGF->Builder;
   assert(E->getBuiltinCallee() == Builtin::BIprintf);
   assert(E->getNumArgs() >= 1); // printf always has at least one arg.
@@ -155,9 +156,15 @@ RValue EmitDevicePrintfCallExpr(const CallExpr *E, CodeGenFunction *CGF,
 
   auto r = packArgsIntoNVPTXFormatBuffer(CGF, Args);
   llvm::Value *BufferPtr = r.first;
+  llvm::Value *Fmt = Args[0].getRValue(*CGF).getScalarVal();
 
-  llvm::SmallVector<llvm::Value *, 3> Vec = {
-      Args[0].getRValue(*CGF).getScalarVal(), BufferPtr};
+  // For OpenCL, the default addrspace of 'format' argument is LangAS::opencl_constant,
+  // however, the 'vprintf' requires it to be unqualified 'ptr' type. Do pointer cast if
+  // it's the case.
+  if (CGM.getContext().getLangOpts().OpenCL)
+    Fmt = Builder.CreatePointerCast(Fmt, llvm::PointerType::getUnqual(Ctx));
+
+  llvm::SmallVector<llvm::Value *, 3> Vec = {Fmt, BufferPtr};
   if (WithSizeArg) {
     // Passing > 32bit of data as a local alloca doesn't work for nvptx or
     // amdgpu

clang/lib/Sema/SemaDecl.cpp

@@ -7008,8 +7008,12 @@ bool Sema::inferObjCARCLifetime(ValueDecl *decl) {
 }
 
 void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
-  if (Decl->getType().hasAddressSpace())
-    return;
+  // Address space is only meaningful for pointer type
+  if (Decl->getType()->isPointerType()) {
+    const PointerType *T = dyn_cast<PointerType>(Decl->getType().getTypePtr());
+    if (T->getPointeeType().hasAddressSpace())
+      return;
+  }
   if (Decl->getType()->isDependentType())
     return;
   if (VarDecl *Var = dyn_cast<VarDecl>(Decl)) {

clang/test/CodeGenOpenCL/test-printf-nvptx.cl

@@ -6,7 +6,7 @@ int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)))
 // NV-LABEL: define dso_local spir_kernel void @test_printf(
 // NV-SAME: ) #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !4 !kernel_arg_access_qual !4 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !4 {
 // NV-NEXT:  entry:
-// NV-NEXT:    [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str) #[[ATTR2:[0-9]+]]
+// NV-NEXT:    [[TMP0:%.*]] = call i32 @vprintf(ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr null)
 // NV-NEXT:    ret void
 //
 __kernel void test_printf() {