diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index f17e4a83305bf..e051cbc648635 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -5710,6 +5710,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, Value *HalfVal = Builder.CreateLoad(Address); return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getFloatTy())); } + case Builtin::BI__builtin_printf: case Builtin::BIprintf: if (getTarget().getTriple().isNVPTX() || getTarget().getTriple().isAMDGCN()) { diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp index e465789a003eb..bd95541647bcf 100644 --- a/clang/lib/CodeGen/CGGPUBuiltin.cpp +++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp @@ -136,7 +136,8 @@ 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->getBuiltinCallee() == Builtin::BIprintf || + E->getBuiltinCallee() == Builtin::BI__builtin_printf); assert(E->getNumArgs() >= 1); // printf always has at least one arg. // Uses the same format as nvptx for the argument packing, but also passes diff --git a/clang/test/CodeGenCUDA/printf-builtin.cu b/clang/test/CodeGenCUDA/printf-builtin.cu new file mode 100644 index 0000000000000..e018d533ed32d --- /dev/null +++ b/clang/test/CodeGenCUDA/printf-builtin.cu @@ -0,0 +1,21 @@ +// REQUIRES: x86-registered-target +// REQUIRES: nvptx-registered-target +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -emit-llvm -disable-llvm-optzns -fno-builtin-printf -fcuda-is-device \ +// RUN: -o - %s | FileCheck %s + +#define __device__ __attribute__((device)) + +extern "C" __device__ int printf(const char *format, ...); + +// CHECK-LABEL: @_Z4foo1v() +__device__ int foo1() { + // CHECK: call i32 @vprintf + // CHECK-NOT: call i32 (ptr, ...) @printf + return __builtin_printf("Hello World\n"); +} + +// CHECK-LABEL: @_Z4foo2v() +__device__ int foo2() { + // CHECK: call i32 (ptr, ...) @printf + return printf("Hello World\n"); +} diff --git a/clang/test/CodeGenHIP/printf-builtin.hip b/clang/test/CodeGenHIP/printf-builtin.hip new file mode 100644 index 0000000000000..df1fbbb6d637a --- /dev/null +++ b/clang/test/CodeGenHIP/printf-builtin.hip @@ -0,0 +1,23 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \ +// RUN: -o - %s | FileCheck --check-prefixes=CHECK,HOSTCALL %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \ +// RUN: -o - %s | FileCheck --check-prefixes=CHECK,BUFFERED %s + +#define __device__ __attribute__((device)) + +extern "C" __device__ int printf(const char *format, ...); + +// CHECK-LABEL: @_Z4foo1v() +__device__ int foo1() { + // HOSTCALL: call i64 @__ockl_printf_begin + // BUFFERED: call ptr addrspace(1) @__printf_alloc + // CHECK-NOT: call i32 (ptr, ...) @printf + return __builtin_printf("Hello World\n"); +} + +// CHECK-LABEL: @_Z4foo2v() +__device__ int foo2() { + // CHECK: call i32 (ptr, ...) @printf + return printf("Hello World\n"); +}