[clang] [clang][AMDGPU][CUDA] Handle __builtin_printf for device printf (PR #68515)

[Mészáros Gergely via cfe-commits]

https://github.com/Maetveis created 
https://github.com/llvm/llvm-project/pull/68515

Previously `__builtin_printf` would result to emitting call to `printf`, even 
though directly calling `printf` was translated.

Ref: #68478

From 193c652ac8a7a387ce99a63cd09b17bdc44f20a1 Mon Sep 17 00:00:00 2001
From: Gergely Meszaros <gerg...@streamhpc.com>
Date: Sun, 8 Oct 2023 09:30:24 +0000
Subject: [PATCH] [clang][AMDGPU][CUDA] Handle __builtin_printf for device
 printf

Previously __builtin_printf would result to emitting call to printf,
even though directly calling printf was translated.

Ref: #68478
---
 clang/lib/CodeGen/CGBuiltin.cpp          |  1 +
 clang/lib/CodeGen/CGGPUBuiltin.cpp       |  3 ++-
 clang/test/CodeGenCUDA/printf-builtin.cu | 20 ++++++++++++++++++++
 clang/test/CodeGenHIP/printf-builtin.hip | 21 +++++++++++++++++++++
 4 files changed, 44 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/CodeGenCUDA/printf-builtin.cu
 create mode 100644 clang/test/CodeGenHIP/printf-builtin.hip

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index bf984861bccb5cc..c16c005787ca778 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5464,6 +5464,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 75fb06de938425d..794be0520163157 100644
--- a/clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -135,7 +135,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 000000000000000..586d00a878ddf89
--- /dev/null
+++ b/clang/test/CodeGenCUDA/printf-builtin.cu
@@ -0,0 +1,20 @@
+// 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-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 000000000000000..76c7d41376c972d
--- /dev/null
+++ b/clang/test/CodeGenHIP/printf-builtin.hip
@@ -0,0 +1,21 @@
+// 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 %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 %s
+
+#define __device__ __attribute__((device))
+
+extern "C" __device__ int printf(const char *format, ...);
+
+// CHECK-LABEL: @_Z4foo1v()
+__device__ int foo1() {
+  // 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");
+}

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits