[OpenMP] Lower printf to __llvm_omp_vprintf

[JonChesterfield]

Summary

Extension of D112504. Lower amdgpu printf to __llvm_omp_vprintf
which takes the same const char*, void* arguments as cuda vprintf and also
passes the size of the void* alloca which will be needed by a non-stub
implementation of __llvm_omp_vprintf for amdgpu.

This removes the amdgpu link error on any printf in a target region in favour
of silently compiling code that doesn't print anything to stdout.

[JonChesterfield]

Makes more sense with D112227 landed, but probably worthwhile just for the old runtime as it removes the link failure from all the tests that use printf.

clang/lib/CodeGen/CGGPUBuiltin.cpp

92

This packArgsIntoNVPTXFormatBuffer helper could/should be factored out as a first patch to make the minimal change to EmitNVPTXDevicePrintfCallExpr clearer in the diff

openmp/libomptarget/DeviceRTL/include/Debug.h

44

the ## could be rolled into some other change, it makes PRINTF("no args") work, at which point we probably don't need/want #define PRINT()

openmp/libomptarget/DeviceRTL/include/Interface.h

351 ↗

(On Diff #382868)

Possibly not a good thing to have in the interface

[JonChesterfield]

• rebase to not require new runtime, stubs now return error

[JonChesterfield]

New runtime failing CI suggests it may take some iteration to land, so reworked this to make sense with or without it.

Doesn't make much difference to nvptx - printf still ultimately resolves to the same vprintf function as before - but a lot of tests that were previously disabled for amdgpu because references to printf were a linker error will now build and run, some of them successfully.

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

192

this should be -1, for 'error'

[JonChesterfield]

Ping

[jdoerfert]

LG, one nit.

clang/lib/CodeGen/CGGPUBuiltin.cpp

250

Suggestion: I feel this could be in a helper to avoid the duplication with the nvptx version. All but the extra argument is the same, no?

openmp/libomptarget/DeviceRTL/src/Debug.cpp

55

Namespace above should include _OMP, be anonymous or the functions should be static.

[JonChesterfield]

Thanks for the review! Couple of helpers to split out and the rebase on main is messy, retesting now

clang/lib/CodeGen/CGGPUBuiltin.cpp

250

Nice spot, yes - the tails of the two functions can fold.

// We don't know how to emit non-scalar varargs.

the block starting with that occurs three times in the files as of this diff, probably also worth splitting out

openmp/libomptarget/DeviceRTL/src/Debug.cpp

55

Sure, will patch. I thought the DeviceRTL had an internalize call as part of the build process which cuts out all the internal symbols but I might be imagining that.

[JonChesterfield]

• rebase, static functions

[JonChesterfield]

• helper for nonscalars

[JonChesterfield]

Updated. Left Builder.CreateCall alone as I've failed to think of a clean way to handle the sometimes present size argument. Tests pass locally and are a credible guess at what will pass the CI bot, but may need to roll back if there's a difference between local and CI

[JonChesterfield]

• fold nvptx and openmp printf into one call
• Drop ReturnValue argument

[JonChesterfield]

Folded EmitOpenMPDevicePrintfCallExpr and EmitNVPTXDevicePrintfCallExpr into asserts plus one line call into a common function. Leaving amdgpu alone in this patch as it's doing something odd with getScalarVal and I don't want to run this through the internal CI stack.

[Conanap]

Hello,

looks like this commit might be causing some failures on PowerPC buildbots, would you be able to take a quick look? Failing test case is Clang :: OpenMP/nvptx_target_printf_codegen.c.

Buildbot link: https://lab.llvm.org/buildbot/#/builders/105/builds/17248/steps/7/logs/FAIL__Clang__nvptx_target_printf_codegen_c

Thanks!

[JonChesterfield]

This passed the openmp buildbot but failed some others, e.g.

clang-armv7-quick/llvm/clang/test/OpenMP/nvptx_target_printf_codegen.c:64:19: 
error: CHECK-64-NEXT: expected string not found in input
// CHECK-64-NEXT: [[TMP6:%.*]] = call i32 @vprintf(i8* [[TMP1]], i8* [[TMP5]])

Reverted while looking into this

Hi @Conanap! Sorry I didn't catch this quicker, you should be green on the next build

[Conanap]

thanks for reverting!

[JonChesterfield]

• ./llvm/utils/update_cc_test_checks.py clang/test/OpenMP/nvptx_target_printf_codegen.c --opt=$HOME/llvm-install/bin/opt

^ that didn't work, updating the diff with the result while trying to remember/derive how this update script is meant to be used

Command Output (stderr):
--
$HOME/llvm-project/clang/test/OpenMP/nvptx_target_printf_codegen.c:49:20: error: CHECK-64-LABEL: expected string not found in input
// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13_worker
                   ^
<stdin>:1:1: note: scanning from here
; ModuleID = '$HOME/llvm-project/clang/test/OpenMP/nvptx_target_printf_codegen.c'
^
<stdin>:21:1: note: possible intended match here
define weak void @__omp_offloading_10303_5e366c_CheckSimple_l13() #0 {
^

I think what has happened is the codegen test case is expecting a _worker subfunction which, after this change, has been inlined. Also, that update_cc_test_checks.py isn't doing what I hoped, e.g. the calls in the test are still to @vprintf and not to @__llvm_oimp_printf.

It's not obviously a problem that the inlining decision has changed though I don't see why this patch would change it. It is a problem that the update script hasn't changed the print name. Deleting everything other than the RUN: statements and re-executing it doesn't update the print function name, though executing the RUN statement in the shell does:
%4 = call i32 @__llvm_omp_vprintf(i8* getelementptr inbounds ([3 x i8], [3 x i8]* @.str2, i64 0, i64 0), i8* %3, i32 4)

[JonChesterfield]

• regen test file using the better './llvm/utils/update_cc_test_checks.py --llvm-bin=/home/amd/llvm-install/bin clang/test/OpenMP/nvptx_target_printf_codegen.c', now passing

[ronlieb]

@JonChesterfield i am seeing the testcase bug50022.cpp occasionally randomly failing in amdgpu bbot
most recently: https://lab.llvm.org/buildbot/#/builders/193/builds/24230

the same bbot in Staging passed on the patch

the test hangs, and after 23 minutes or so, the bbot terminates.

clang/lib/CodeGen/CGBuiltin.cpp

case Builtin::BI__builtin_load_half: {
     return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getDoubleTy()));
   }
   case Builtin::BI__builtin_load_halff: {
     Address Address = EmitPointerWithAlignment(E->getArg(0));
     Value *HalfVal = Builder.CreateLoad(Address);
     return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getFloatTy()));
   }
   case Builtin::BIprintf:
+    if (getTarget().getTriple().isNVPTX() ||
+        getTarget().getTriple().isAMDGCN()) {
+      if (getLangOpts().OpenMPIsDevice)
+        return EmitOpenMPDevicePrintfCallExpr(E);
       if (getTarget().getTriple().isNVPTX())
-      return EmitNVPTXDevicePrintfCallExpr(E, ReturnValue);
-    if (getTarget().getTriple().getArch() == Triple::amdgcn &&
-        getLangOpts().HIP)
-      return EmitAMDGPUDevicePrintfCallExpr(E, ReturnValue);
+        return EmitNVPTXDevicePrintfCallExpr(E);
+      if (getTarget().getTriple().isAMDGCN() && getLangOpts().HIP)
+        return EmitAMDGPUDevicePrintfCallExpr(E);
+    }
+
     break;
   case Builtin::BI__builtin_canonicalize:
   case Builtin::BI__builtin_canonicalizef:
   case Builtin::BI__builtin_canonicalizef16:
   case Builtin::BI__builtin_canonicalizel:
     return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::canonicalize));
 
   case Builtin::BI__builtin_thread_pointer: {

clang/lib/CodeGen/CGGPUBuiltin.cpp

 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/Instruction.h"
 #include "llvm/Support/MathExtras.h"
 #include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
 
 using namespace clang;
 using namespace CodeGen;
 
-static llvm::Function *GetVprintfDeclaration(llvm::Module &M) {
+namespace {
+llvm::Function *GetVprintfDeclaration(llvm::Module &M) {
   llvm::Type *ArgTypes[] = {llvm::Type::getInt8PtrTy(M.getContext()),
                             llvm::Type::getInt8PtrTy(M.getContext())};
   llvm::FunctionType *VprintfFuncType = llvm::FunctionType::get(
       llvm::Type::getInt32Ty(M.getContext()), ArgTypes, false);
 
-  if (auto* F = M.getFunction("vprintf")) {
+  if (auto *F = M.getFunction("vprintf")) {
     // Our CUDA system header declares vprintf with the right signature, so
     // nobody else should have been able to declare vprintf with a bogus
     // signature.
     assert(F->getFunctionType() == VprintfFuncType);
     return F;
   }
 
   // vprintf doesn't already exist; create a declaration and insert it into the
   // module.
   return llvm::Function::Create(
       VprintfFuncType, llvm::GlobalVariable::ExternalLinkage, "vprintf", &M);
 }
 
+llvm::Function *GetOpenMPVprintfDeclaration(CodeGenModule &CGM) {
+  const char *Name = "__llvm_omp_vprintf";
+  llvm::Module &M = CGM.getModule();
+  llvm::Type *ArgTypes[] = {llvm::Type::getInt8PtrTy(M.getContext()),
+                            llvm::Type::getInt8PtrTy(M.getContext()),
+                            llvm::Type::getInt32Ty(M.getContext())};
+  llvm::FunctionType *VprintfFuncType = llvm::FunctionType::get(
+      llvm::Type::getInt32Ty(M.getContext()), ArgTypes, false);
+
+  if (auto *F = M.getFunction(Name)) {
+    if (F->getFunctionType() != VprintfFuncType) {
+      CGM.Error(SourceLocation(),
+                "Invalid type declaration for __llvm_omp_vprintf");
+      return nullptr;
+    }
+    return F;
+  }
+
+  return llvm::Function::Create(
+      VprintfFuncType, llvm::GlobalVariable::ExternalLinkage, Name, &M);
+}
+
 // Transforms a call to printf into a call to the NVPTX vprintf syscall (which
 // isn't particularly special; it's invoked just like a regular function).
 // vprintf takes two args: A format string, and a pointer to a buffer containing
 // the varargs.
 //
 // For example, the call
 //
 //   printf("format string", arg1, arg2, arg3);

 //   *(Tmp*)buf = {a1, a2, a3};
 //   vprintf("format string", buf);
 //
 // buf is aligned to the max of {alignof(Arg1), ...}.  Furthermore, each of the
 // args is itself aligned to its preferred alignment.
 //
 // Note that by the time this function runs, E's args have already undergone the
 // standard C vararg promotion (short -> int, float -> double, etc.).
 
-namespace {
-llvm::Value *packArgsIntoNVPTXFormatBuffer(CodeGenFunction *CGF,
-                                           const CallArgList &Args) {
+std::pair<llvm::Value *, llvm::TypeSize>
+packArgsIntoNVPTXFormatBuffer(CodeGenFunction *CGF, const CallArgList &Args) {
   const llvm::DataLayout &DL = CGF->CGM.getDataLayout();
   llvm::LLVMContext &Ctx = CGF->CGM.getLLVMContext();
   CGBuilderTy &Builder = CGF->Builder;
 
   // Construct and fill the args buffer that we'll pass to vprintf.
   if (Args.size() <= 1) {
-    // If there are no args, pass a null pointer to vprintf.
-    return llvm::ConstantPointerNull::get(llvm::Type::getInt8PtrTy(Ctx));
+    // If there are no args, pass a null pointer and size 0
+    llvm::Value * BufferPtr = llvm::ConstantPointerNull::get(llvm::Type::getInt8PtrTy(Ctx));
+    return {BufferPtr, llvm::TypeSize::Fixed(0)};
   } else {
     llvm::SmallVector<llvm::Type *, 8> ArgTypes;
     for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I)
       ArgTypes.push_back(Args[I].getRValue(*CGF).getScalarVal()->getType());
 
     // Using llvm::StructType is correct only because printf doesn't accept
     // aggregates.  If we had to handle aggregates here, we'd have to manually
     // compute the offsets within the alloca -- we wouldn't be able to assume
     // that the alignment of the llvm type was the same as the alignment of the
     // clang type.
     llvm::Type *AllocaTy = llvm::StructType::create(ArgTypes, "printf_args");
     llvm::Value *Alloca = CGF->CreateTempAlloca(AllocaTy);
 
     for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I) {
       llvm::Value *P = Builder.CreateStructGEP(AllocaTy, Alloca, I - 1);
       llvm::Value *Arg = Args[I].getRValue(*CGF).getScalarVal();
       Builder.CreateAlignedStore(Arg, P, DL.getPrefTypeAlign(Arg->getType()));
     }
-    return Builder.CreatePointerCast(Alloca, llvm::Type::getInt8PtrTy(Ctx));
+    llvm::Value *BufferPtr =
+        Builder.CreatePointerCast(Alloca, llvm::Type::getInt8PtrTy(Ctx));
+    return {BufferPtr, DL.getTypeAllocSize(AllocaTy)};
   }
 }
-} // namespace
 
-RValue
-CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E,
-                                               ReturnValueSlot ReturnValue) {
-  assert(getTarget().getTriple().isNVPTX());
+bool containsNonScalarVarargs(CodeGenFunction *CGF, CallArgList Args) {
+  return llvm::any_of(llvm::drop_begin(Args), [&](const CallArg &A) {
+    return !A.getRValue(*CGF).isScalar();
+  });
+}
+
+RValue EmitDevicePrintfCallExpr(const CallExpr *E, CodeGenFunction *CGF,
+                                llvm::Function *Decl, bool WithSizeArg) {
+  CodeGenModule &CGM = CGF->CGM;
+  CGBuilderTy &Builder = CGF->Builder;
   assert(E->getBuiltinCallee() == Builtin::BIprintf);
   assert(E->getNumArgs() >= 1); // printf always has at least one arg.
 
+  // Uses the same format as nvptx for the argument packing, but also passes
+  // an i32 for the total size of the passed pointer
   CallArgList Args;
-  EmitCallArgs(Args,
+  CGF->EmitCallArgs(Args,
                     E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
                     E->arguments(), E->getDirectCallee(),
                     /* ParamsToSkip = */ 0);
 
   // We don't know how to emit non-scalar varargs.
-  if (llvm::any_of(llvm::drop_begin(Args), [&](const CallArg &A) {
-        return !A.getRValue(*this).isScalar();
-      })) {
+  if (containsNonScalarVarargs(CGF, Args)) {
     CGM.ErrorUnsupported(E, "non-scalar arg to printf");
-    return RValue::get(llvm::ConstantInt::get(IntTy, 0));
+    return RValue::get(llvm::ConstantInt::get(CGF->IntTy, 0));
   }
 
-  llvm::Value *BufferPtr = packArgsIntoNVPTXFormatBuffer(this, Args);
+  auto r = packArgsIntoNVPTXFormatBuffer(CGF, Args);
+  llvm::Value *BufferPtr = r.first;
+
+  llvm::SmallVector<llvm::Value *, 3> Vec = {
+      Args[0].getRValue(*CGF).getScalarVal(), BufferPtr};
+  if (WithSizeArg) {
+    // Passing > 32bit of data as a local alloca doesn't work for nvptx or
+    // amdgpu
+    llvm::Constant *Size =
+        llvm::ConstantInt::get(llvm::Type::getInt32Ty(CGM.getLLVMContext()),
+                               static_cast<uint32_t>(r.second.getFixedSize()));
 
-  // Invoke vprintf and return.
-  llvm::Function* VprintfFunc = GetVprintfDeclaration(CGM.getModule());
-  return RValue::get(Builder.CreateCall(
-      VprintfFunc, {Args[0].getRValue(*this).getScalarVal(), BufferPtr}));
+    Vec.push_back(Size);
+  }
+  return RValue::get(Builder.CreateCall(Decl, Vec));
+}
+} // namespace
+
+RValue CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E) {
+  assert(getTarget().getTriple().isNVPTX());
+  return EmitDevicePrintfCallExpr(
+      E, this, GetVprintfDeclaration(CGM.getModule()), false);
 }
 
-RValue
-CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E,
-                                                ReturnValueSlot ReturnValue) {
+RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
   assert(getTarget().getTriple().getArch() == llvm::Triple::amdgcn);
   assert(E->getBuiltinCallee() == Builtin::BIprintf ||
          E->getBuiltinCallee() == Builtin::BI__builtin_printf);
   assert(E->getNumArgs() >= 1); // printf always has at least one arg.
 
   CallArgList CallArgs;
   EmitCallArgs(CallArgs,
                E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
   }
 
   llvm::IRBuilder<> IRB(Builder.GetInsertBlock(), Builder.GetInsertPoint());
   IRB.SetCurrentDebugLocation(Builder.getCurrentDebugLocation());
   auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args);
   Builder.SetInsertPoint(IRB.GetInsertBlock(), IRB.GetInsertPoint());
   return RValue::get(Printf);
 }
+
+RValue CodeGenFunction::EmitOpenMPDevicePrintfCallExpr(const CallExpr *E) {
+  assert(getTarget().getTriple().isNVPTX() ||
+         getTarget().getTriple().isAMDGCN());
+  return EmitDevicePrintfCallExpr(E, this, GetOpenMPVprintfDeclaration(CGM),
+                                  true);
+}

clang/lib/CodeGen/CodeGenFunction.h

public:
   RValue EmitCXXOperatorMemberCallExpr(const CXXOperatorCallExpr *E,
                                        const CXXMethodDecl *MD,
                                        ReturnValueSlot ReturnValue);
   RValue EmitCXXPseudoDestructorExpr(const CXXPseudoDestructorExpr *E);
 
   RValue EmitCUDAKernelCallExpr(const CUDAKernelCallExpr *E,
                                 ReturnValueSlot ReturnValue);
 
-  RValue EmitNVPTXDevicePrintfCallExpr(const CallExpr *E,
-                                       ReturnValueSlot ReturnValue);
-  RValue EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E,
-                                        ReturnValueSlot ReturnValue);
+  RValue EmitNVPTXDevicePrintfCallExpr(const CallExpr *E);
+  RValue EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E);
+  RValue EmitOpenMPDevicePrintfCallExpr(const CallExpr *E);
 
   RValue EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
                          const CallExpr *E, ReturnValueSlot ReturnValue);
 
   RValue emitRotate(const CallExpr *E, bool IsRotateRight);
 
   /// Emit IR for __builtin_os_log_format.
   RValue emitBuiltinOSLogFormat(const CallExpr &E);

clang/test/OpenMP/nvptx_target_printf_codegen.c

 // CHECK-64-NEXT:    [[TMP1:%.*]] = load i8*, i8** [[FMT]], align 8
 // CHECK-64-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 0
 // CHECK-64-NEXT:    store i32 1, i32* [[TMP2]], align 4
 // CHECK-64-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 1
 // CHECK-64-NEXT:    store i64 2, i64* [[TMP3]], align 8
 // CHECK-64-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 2
 // CHECK-64-NEXT:    store double 3.000000e+00, double* [[TMP4]], align 8
 // CHECK-64-NEXT:    [[TMP5:%.*]] = bitcast %printf_args* [[TMP]] to i8*
-// CHECK-64-NEXT:    [[TMP6:%.*]] = call i32 @vprintf(i8* [[TMP1]], i8* [[TMP5]])
+// CHECK-64-NEXT:    [[TMP6:%.*]] = call i32 @__llvm_omp_vprintf(i8* [[TMP1]], i8* [[TMP5]], i32 24)
 // CHECK-64-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-64-NEXT:    ret void
 // CHECK-64:       worker.exit:
 // CHECK-64-NEXT:    ret void
 //
 //
 // CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25
 // CHECK-64-SAME: () #[[ATTR0]] {
 // CHECK-64-NEXT:  entry:
 // CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 true, i1 true)
 // CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 // CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 // CHECK-64:       user_code.entry:
-// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @vprintf(i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str1, i64 0, i64 0), i8* null)
+// CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__llvm_omp_vprintf(i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str1, i64 0, i64 0), i8* null, i32 0)
 // CHECK-64-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-64-NEXT:    ret void
 // CHECK-64:       worker.exit:
 // CHECK-64-NEXT:    ret void
 //
 //
 // CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36
 // CHECK-64-SAME: (i64 [[FOO:%.*]]) #[[ATTR0]] {
 // CHECK-64-NEXT:  entry:
 // CHECK-64-NEXT:    [[FOO_ADDR:%.*]] = alloca i64, align 8
 // CHECK-64-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
 // CHECK-64-NEXT:    store i64 [[FOO]], i64* [[FOO_ADDR]], align 8
 // CHECK-64-NEXT:    [[CONV:%.*]] = bitcast i64* [[FOO_ADDR]] to i32*
 // CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 true, i1 true)
 // CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 // CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 // CHECK-64:       user_code.entry:
 // CHECK-64-NEXT:    [[TMP1:%.*]] = load i32, i32* [[CONV]], align 8
 // CHECK-64-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP1]], 0
 // CHECK-64-NEXT:    br i1 [[TOBOOL]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
 // CHECK-64:       if.then:
 // CHECK-64-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS_0]], %printf_args.0* [[TMP]], i32 0, i32 0
 // CHECK-64-NEXT:    store i32 42, i32* [[TMP2]], align 4
 // CHECK-64-NEXT:    [[TMP3:%.*]] = bitcast %printf_args.0* [[TMP]] to i8*
-// CHECK-64-NEXT:    [[TMP4:%.*]] = call i32 @vprintf(i8* getelementptr inbounds ([3 x i8], [3 x i8]* @.str2, i64 0, i64 0), i8* [[TMP3]])
+// CHECK-64-NEXT:    [[TMP4:%.*]] = call i32 @__llvm_omp_vprintf(i8* getelementptr inbounds ([3 x i8], [3 x i8]* @.str2, i64 0, i64 0), i8* [[TMP3]], i32 4)
 // CHECK-64-NEXT:    br label [[IF_END]]
 // CHECK-64:       worker.exit:
 // CHECK-64-NEXT:    ret void
 // CHECK-64:       if.end:
 // CHECK-64-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-64-NEXT:    ret void
 //
 //

 // CHECK-32-NEXT:    [[TMP1:%.*]] = load i8*, i8** [[FMT]], align 4
 // CHECK-32-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 0
 // CHECK-32-NEXT:    store i32 1, i32* [[TMP2]], align 4
 // CHECK-32-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 1
 // CHECK-32-NEXT:    store i64 2, i64* [[TMP3]], align 8
 // CHECK-32-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], %printf_args* [[TMP]], i32 0, i32 2
 // CHECK-32-NEXT:    store double 3.000000e+00, double* [[TMP4]], align 8
 // CHECK-32-NEXT:    [[TMP5:%.*]] = bitcast %printf_args* [[TMP]] to i8*
-// CHECK-32-NEXT:    [[TMP6:%.*]] = call i32 @vprintf(i8* [[TMP1]], i8* [[TMP5]])
+// CHECK-32-NEXT:    [[TMP6:%.*]] = call i32 @__llvm_omp_vprintf(i8* [[TMP1]], i8* [[TMP5]], i32 24)
 // CHECK-32-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-32-NEXT:    ret void
 // CHECK-32:       worker.exit:
 // CHECK-32-NEXT:    ret void
 //
 //
 // CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25
 // CHECK-32-SAME: () #[[ATTR0]] {
 // CHECK-32-NEXT:  entry:
 // CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 true, i1 true)
 // CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 // CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 // CHECK-32:       user_code.entry:
-// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @vprintf(i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str1, i32 0, i32 0), i8* null)
+// CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__llvm_omp_vprintf(i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str1, i32 0, i32 0), i8* null, i32 0)
 // CHECK-32-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-32-NEXT:    ret void
 // CHECK-32:       worker.exit:
 // CHECK-32-NEXT:    ret void
 //
 //
 // CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36
 // CHECK-32-SAME: (i32 [[FOO:%.*]]) #[[ATTR0]] {
 // CHECK-32-NEXT:  entry:
 // CHECK-32-NEXT:    [[FOO_ADDR:%.*]] = alloca i32, align 4
 // CHECK-32-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
 // CHECK-32-NEXT:    store i32 [[FOO]], i32* [[FOO_ADDR]], align 4
 // CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 1, i1 true, i1 true)
 // CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 // CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 // CHECK-32:       user_code.entry:
 // CHECK-32-NEXT:    [[TMP1:%.*]] = load i32, i32* [[FOO_ADDR]], align 4
 // CHECK-32-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP1]], 0
 // CHECK-32-NEXT:    br i1 [[TOBOOL]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
 // CHECK-32:       if.then:
 // CHECK-32-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS_0]], %printf_args.0* [[TMP]], i32 0, i32 0
 // CHECK-32-NEXT:    store i32 42, i32* [[TMP2]], align 4
 // CHECK-32-NEXT:    [[TMP3:%.*]] = bitcast %printf_args.0* [[TMP]] to i8*
-// CHECK-32-NEXT:    [[TMP4:%.*]] = call i32 @vprintf(i8* getelementptr inbounds ([3 x i8], [3 x i8]* @.str2, i32 0, i32 0), i8* [[TMP3]])
+// CHECK-32-NEXT:    [[TMP4:%.*]] = call i32 @__llvm_omp_vprintf(i8* getelementptr inbounds ([3 x i8], [3 x i8]* @.str2, i32 0, i32 0), i8* [[TMP3]], i32 4)
 // CHECK-32-NEXT:    br label [[IF_END]]
 // CHECK-32:       worker.exit:
 // CHECK-32-NEXT:    ret void
 // CHECK-32:       if.end:
 // CHECK-32-NEXT:    call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
 // CHECK-32-NEXT:    ret void
 //

openmp/libomptarget/DeviceRTL/include/Debug.h

if (config::isDebugMode(config::DebugKind::Assertion) && !(expr))          \
       __assert_fail(#expr, __FILE__, __LINE__, __PRETTY_FUNCTION__);           \
     else                                                                       \
       __assert_assume(expr);                                                   \
   }
 
 ///}
 
 /// Print
-/// TODO: For now we have to use macros to guard the code because Clang lowers
-/// `printf` to different function calls on NVPTX and AMDGCN platforms, and it
-/// doesn't work for AMDGCN. After it can work on AMDGCN, we will remove the
-/// macro.
+/// printf() calls are rewritten by CGGPUBuiltin to __llvm_omp_vprintf
 /// {
 
-#ifndef __AMDGCN__
 extern "C" {
 int printf(const char *format, ...);
 }
 
-#define PRINTF(fmt, ...) (void)printf(fmt, __VA_ARGS__);
+#define PRINTF(fmt, ...) (void)printf(fmt, ##__VA_ARGS__);
 #define PRINT(str) PRINTF("%s", str)
-#else
-#define PRINTF(fmt, ...)
-#define PRINT(str)
-#endif
 
 ///}
 
 /// Enter a debugging scope for performing function traces. Enabled with
 /// FunctionTracting set in the debug kind.
 #define FunctionTracingRAII()                                                  \
   DebugEntryRAII Entry(__FILE__, __LINE__, __PRETTY_FUNCTION__);
 

openmp/libomptarget/DeviceRTL/src/Debug.cpp

 void __assert_assume(bool condition) { __builtin_assume(condition); }
 
 void __assert_fail(const char *assertion, const char *file, unsigned line,
                    const char *function) {
   PRINTF("%s:%u: %s: Assertion `%s' failed.\n", file, line, function,
          assertion);
   __builtin_trap();
 }
+
+#pragma omp begin declare variant match(                                       \
+    device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
+int32_t vprintf(const char *, void *);
+namespace impl {
+static int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
+  return vprintf(Format, Arguments);
+}
+} // namespace impl
+#pragma omp end declare variant
+
+// We do not have a vprintf implementation for AMD GPU yet so we use a stub.
+#pragma omp begin declare variant match(device = {arch(amdgcn)})
+namespace impl {
+static int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
+  return -1;
+}
+} // namespace impl
+#pragma omp end declare variant
+
+int32_t __llvm_omp_vprintf(const char *Format, void *Arguments, uint32_t Size) {
+  return impl::omp_vprintf(Format, Arguments, Size);
+}
 }
 
 /// Current indentation level for the function trace. Only accessed by thread 0.
 static uint32_t Level = 0;
 #pragma omp allocate(Level) allocator(omp_pteam_mem_alloc)
 
 DebugEntryRAII::DebugEntryRAII(const char *File, const unsigned Line,
                                const char *Function) {
   if (config::isDebugMode(config::DebugKind::FunctionTracing) &&

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

 // Stub implementations
 // Weak to allow overriding by local versions while comparing different
 // potential implementations
 __attribute__((weak)) EXTERN void *__kmpc_impl_malloc(size_t) {
   return nullptr;
 }
 __attribute__((weak)) EXTERN void __kmpc_impl_free(void *) {}
 
+EXTERN
+int32_t __llvm_omp_vprintf(const char *Format, void *Arguments, uint32_t) {
+  return -1;
+}
+
 EXTERN void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));
   hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
 }
 
 EXTERN uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   return (((uint64_t)hi) << 32) | (uint64_t)lo;
 }

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

 
 EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock) {
   return __kmpc_atomic_add(lock, 0u);
 }
 
 extern "C" {
 void *malloc(size_t);
 void free(void *);
+int32_t vprintf(const char *, void *);
 }
 
 EXTERN void *__kmpc_impl_malloc(size_t x) { return malloc(x); }
 EXTERN void __kmpc_impl_free(void *x) { free(x); }
 
+EXTERN int32_t __llvm_omp_vprintf(const char *Format, void *Arguments,
+                                  uint32_t) {
+  return vprintf(Format, Arguments);
+}
+
 #pragma omp end declare target

openmp/libomptarget/test/mapping/data_member_ref.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// amdgcn does not have printf definition
+// Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 
 struct View {
   int Data;
 };

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// amdgcn does not have printf definition
+// Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;

openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// amdgcn does not have printf definition
+// Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;

openmp/libomptarget/test/mapping/lambda_by_value.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// amdgcn does not have printf definition
+// Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <stdint.h>
 
 // CHECK: before: [[V1:111]] [[V2:222]] [[PX:0x[^ ]+]] [[PY:0x[^ ]+]]
 // CHECK: lambda: [[V1]] [[V2]] [[PX_TGT:0x[^ ]+]] 0x{{.*}}

openmp/libomptarget/test/mapping/ompx_hold/struct.c

 // RUN: %libomptarget-compile-generic -fopenmp-extensions
 // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
 
-// amdgcn does not have printf definition
+// Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>
 
 #define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
   printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \

openmp/libomptarget/test/mapping/ptr_and_obj_motion.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
-// amdgcn does not have printf definition
-// XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
-
 #include <stdio.h>
 
 typedef struct {
   double *dataptr;
   int dummy1;
   int dummy2;
 } DV;
 

openmp/libomptarget/test/mapping/reduction_implicit_map.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// amdgcn does not have printf definition
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
-
 #include <stdio.h>
 
 void sum(int* input, int size, int* output)
 {
 #pragma omp target teams distribute parallel for reduction(+:output[0]) \
                                                  map(to:input[0:size])
   for (int i = 0; i < size; i++)
     output[0] += input[i];

openmp/libomptarget/test/offloading/bug49021.cpp

 // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
 
-// Wrong results on amdgcn
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
+// Wrong results on amdgpu
+// XFAIL: amdgcn-amd-amdhsa
 
 #include <iostream>
 
 template <typename T> int test_map() {
   std::cout << "map(complex<>)" << std::endl;
   T a(0.2), a_check;
 #pragma omp target map(from : a_check)
   { a_check = a; }

openmp/libomptarget/test/offloading/bug50022.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
 
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
-
 #include <cassert>
 #include <iostream>
 #include <stdexcept>
 
 int main(int argc, char *argv[]) {
   int a = 0;
   std::cout << "outside a = " << a << " addr " << &a << std::endl;
 #pragma omp target map(tofrom : a) depend(out : a) nowait

openmp/libomptarget/test/offloading/host_as_target.c

 // Check that specifying device as omp_get_initial_device():
 // - Doesn't cause the runtime to fail.
 // - Offloads code to the host.
 // - Doesn't transfer data.  In this case, just check that neither host data nor
 //   default device data are affected by the specified transfers.
 // - Works whether it's specified directly or as the default device.
 
 // RUN: %libomptarget-compile-run-and-check-generic
 
-// amdgcn does not have printf definition
+// amdgpu does not have a working printf definition
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 static void check(char *X, int Dev) {
   printf("  host X = %c\n", *X);

openmp/libomptarget/test/unified_shared_memory/api.c

 // RUN: %libomptarget-compile-run-and-check-generic
 // XFAIL: nvptx64-nvidia-cuda
 // XFAIL: nvptx64-nvidia-cuda-newRTL
 
-// Fails on amdgcn with error: GPU Memory Error
+// Fails on amdgpu with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL

openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
 
-// Fails on amdgcn with error: GPU Memory Error
+// Fails on amdgpu with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 

openmp/libomptarget/test/unified_shared_memory/close_modifier.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
 
-// amdgcn does not have printf definition
-// XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
+// amdgpu runtime crash
+// UNSUPPORTED: amdgcn-amd-amdhsa
+
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 
 #define N 1024
 

openmp/libomptarget/test/unified_shared_memory/shared_update.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 
-// amdgcn does not have printf definition
-// XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
+// amdgpu runtime crash
+// UNSUPPORTED: amdgcn-amd-amdhsa
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);