[libc-commits] [libc] dbb8b7a - Reapply "[OpenMP][libc] Remove special handling for OpenMP printf (#98940)"
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Fri Jul 26 15:22:02 PDT 2024
Author: Joseph Huber
Date: 2024-07-26T17:21:56-05:00
New Revision: dbb8b7a0f4eea1aa333cec9a38aa6eb7ecf6c1dc
URL: https://github.com/llvm/llvm-project/commit/dbb8b7a0f4eea1aa333cec9a38aa6eb7ecf6c1dc
DIFF: https://github.com/llvm/llvm-project/commit/dbb8b7a0f4eea1aa333cec9a38aa6eb7ecf6c1dc.diff
LOG: Reapply "[OpenMP][libc] Remove special handling for OpenMP printf (#98940)"
This reverts commit fea5914c926e2f013a8b5e27eaa74c7047fb2c71.
Added:
Modified:
clang/lib/CodeGen/CGBuiltin.cpp
clang/lib/CodeGen/CGGPUBuiltin.cpp
clang/lib/CodeGen/CodeGenFunction.h
libc/config/gpu/entrypoints.txt
libc/spec/gpu_ext.td
libc/src/gpu/CMakeLists.txt
llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
offload/DeviceRTL/include/LibC.h
offload/DeviceRTL/src/LibC.cpp
Removed:
clang/test/OpenMP/nvptx_target_printf_codegen.c
libc/src/gpu/rpc_fprintf.cpp
libc/src/gpu/rpc_fprintf.h
################################################################################
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 0c4d0efb70ea5..f0651c280ff95 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5986,8 +5986,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
getTarget().getTriple().isAMDGCN() ||
(getTarget().getTriple().isSPIRV() &&
getTarget().getTriple().getVendor() == Triple::VendorType::AMD)) {
- if (getLangOpts().OpenMPIsTargetDevice)
- return EmitOpenMPDevicePrintfCallExpr(E);
if (getTarget().getTriple().isNVPTX())
return EmitNVPTXDevicePrintfCallExpr(E);
if ((getTarget().getTriple().isAMDGCN() ||
diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp
index b2340732afeb5..84adf29e8db87 100644
--- a/clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -42,28 +42,6 @@ llvm::Function *GetVprintfDeclaration(llvm::Module &M) {
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::PointerType::getUnqual(M.getContext()),
- llvm::PointerType::getUnqual(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
@@ -213,10 +191,3 @@ RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
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);
-}
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index bd62c65d8cce6..89cc819c43bb5 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4536,7 +4536,6 @@ class CodeGenFunction : public CodeGenTypeCache {
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);
diff --git a/clang/test/OpenMP/nvptx_target_printf_codegen.c b/clang/test/OpenMP/nvptx_target_printf_codegen.c
deleted file mode 100644
index f53daf65205c9..0000000000000
--- a/clang/test/OpenMP/nvptx_target_printf_codegen.c
+++ /dev/null
@@ -1,179 +0,0 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
-// Test target codegen - host bc file has to be created first.
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-64
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK-32
-// expected-no-diagnostics
-extern int printf(const char *, ...);
-
-
-// Check a simple call to printf end-to-end.
-int CheckSimple(void) {
-#pragma omp target
- {
- // printf in master-only basic block.
- const char* fmt = "%d %lld %f";
-
- printf(fmt, 1, 2ll, 3.0);
- }
-
- return 0;
-}
-
-void CheckNoArgs(void) {
-#pragma omp target
- {
- // printf in master-only basic block.
- printf("hello, world!");
- }
-}
-
-// Check that printf's alloca happens in the entry block, not inside the if
-// statement.
-int foo;
-void CheckAllocaIsInEntryBlock(void) {
-#pragma omp target
- {
- if (foo) {
- printf("%d", 42);
- }
- }
-}
-// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13
-// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
-// CHECK-64-NEXT: entry:
-// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-64-NEXT: [[FMT:%.*]] = alloca ptr, align 8
-// CHECK-64-NEXT: [[TMP:%.*]] = alloca [[PRINTF_ARGS:%.*]], align 8
-// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
-// CHECK-64-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13_kernel_environment, ptr [[DYN_PTR]])
-// 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: store ptr @.str, ptr [[FMT]], align 8
-// CHECK-64-NEXT: [[TMP1:%.*]] = load ptr, ptr [[FMT]], align 8
-// CHECK-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 0
-// CHECK-64-NEXT: store i32 1, ptr [[TMP2]], align 4
-// CHECK-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 1
-// CHECK-64-NEXT: store i64 2, ptr [[TMP3]], align 8
-// CHECK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 2
-// CHECK-64-NEXT: store double 3.000000e+00, ptr [[TMP4]], align 8
-// CHECK-64-NEXT: [[TMP5:%.*]] = call i32 @__llvm_omp_vprintf(ptr [[TMP1]], ptr [[TMP]], i32 24)
-// CHECK-64-NEXT: call void @__kmpc_target_deinit()
-// 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: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
-// CHECK-64-NEXT: entry:
-// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
-// CHECK-64-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25_kernel_environment, ptr [[DYN_PTR]])
-// 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 @__llvm_omp_vprintf(ptr @.str1, ptr null, i32 0)
-// CHECK-64-NEXT: call void @__kmpc_target_deinit()
-// 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: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[FOO:%.*]]) #[[ATTR0]] {
-// CHECK-64-NEXT: entry:
-// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-64-NEXT: [[FOO_ADDR:%.*]] = alloca i64, align 8
-// CHECK-64-NEXT: [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
-// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
-// CHECK-64-NEXT: store i64 [[FOO]], ptr [[FOO_ADDR]], align 8
-// CHECK-64-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36_kernel_environment, ptr [[DYN_PTR]])
-// 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, ptr [[FOO_ADDR]], align 4
-// 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]], ptr [[TMP]], i32 0, i32 0
-// CHECK-64-NEXT: store i32 42, ptr [[TMP2]], align 4
-// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str2, ptr [[TMP]], 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()
-// CHECK-64-NEXT: ret void
-//
-//
-// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13
-// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
-// CHECK-32-NEXT: entry:
-// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
-// CHECK-32-NEXT: [[FMT:%.*]] = alloca ptr, align 4
-// CHECK-32-NEXT: [[TMP:%.*]] = alloca [[PRINTF_ARGS:%.*]], align 8
-// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
-// CHECK-32-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13_kernel_environment, ptr [[DYN_PTR]])
-// 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: store ptr @.str, ptr [[FMT]], align 4
-// CHECK-32-NEXT: [[TMP1:%.*]] = load ptr, ptr [[FMT]], align 4
-// CHECK-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 0
-// CHECK-32-NEXT: store i32 1, ptr [[TMP2]], align 4
-// CHECK-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 1
-// CHECK-32-NEXT: store i64 2, ptr [[TMP3]], align 8
-// CHECK-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 2
-// CHECK-32-NEXT: store double 3.000000e+00, ptr [[TMP4]], align 8
-// CHECK-32-NEXT: [[TMP5:%.*]] = call i32 @__llvm_omp_vprintf(ptr [[TMP1]], ptr [[TMP]], i32 24)
-// CHECK-32-NEXT: call void @__kmpc_target_deinit()
-// 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: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
-// CHECK-32-NEXT: entry:
-// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
-// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
-// CHECK-32-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25_kernel_environment, ptr [[DYN_PTR]])
-// 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 @__llvm_omp_vprintf(ptr @.str1, ptr null, i32 0)
-// CHECK-32-NEXT: call void @__kmpc_target_deinit()
-// 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: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[FOO:%.*]]) #[[ATTR0]] {
-// CHECK-32-NEXT: entry:
-// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
-// CHECK-32-NEXT: [[FOO_ADDR:%.*]] = alloca i32, align 4
-// CHECK-32-NEXT: [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
-// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
-// CHECK-32-NEXT: store i32 [[FOO]], ptr [[FOO_ADDR]], align 4
-// CHECK-32-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36_kernel_environment, ptr [[DYN_PTR]])
-// 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, ptr [[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]], ptr [[TMP]], i32 0, i32 0
-// CHECK-32-NEXT: store i32 42, ptr [[TMP2]], align 4
-// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str2, ptr [[TMP]], 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()
-// CHECK-32-NEXT: ret void
-//
diff --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
index df7aa9e319624..157f6f8af00a9 100644
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -226,7 +226,6 @@ set(TARGET_LIBC_ENTRYPOINTS
# gpu/rpc.h entrypoints
libc.src.gpu.rpc_host_call
- libc.src.gpu.rpc_fprintf
)
set(TARGET_LIBM_ENTRYPOINTS
diff --git a/libc/spec/gpu_ext.td b/libc/spec/gpu_ext.td
index 5400e0afa7564..dce81ff778620 100644
--- a/libc/spec/gpu_ext.td
+++ b/libc/spec/gpu_ext.td
@@ -10,14 +10,6 @@ def GPUExtensions : StandardSpec<"GPUExtensions"> {
RetValSpec<VoidType>,
[ArgSpec<VoidPtr>, ArgSpec<VoidPtr>, ArgSpec<SizeTType>]
>,
- FunctionSpec<
- "rpc_fprintf",
- RetValSpec<IntType>,
- [ArgSpec<FILERestrictedPtr>,
- ArgSpec<ConstCharRestrictedPtr>,
- ArgSpec<VoidPtr>,
- ArgSpec<SizeTType>]
- >,
]
>;
let Headers = [
diff --git a/libc/src/gpu/CMakeLists.txt b/libc/src/gpu/CMakeLists.txt
index 4508abea7a888..e20228516b511 100644
--- a/libc/src/gpu/CMakeLists.txt
+++ b/libc/src/gpu/CMakeLists.txt
@@ -8,15 +8,3 @@ add_entrypoint_object(
libc.src.__support.RPC.rpc_client
libc.src.__support.GPU.utils
)
-
-add_entrypoint_object(
- rpc_fprintf
- SRCS
- rpc_fprintf.cpp
- HDRS
- rpc_fprintf.h
- DEPENDS
- libc.src.stdio.gpu.gpu_file
- libc.src.__support.RPC.rpc_client
- libc.src.__support.GPU.utils
-)
diff --git a/libc/src/gpu/rpc_fprintf.cpp b/libc/src/gpu/rpc_fprintf.cpp
deleted file mode 100644
index 70056daa25e2e..0000000000000
--- a/libc/src/gpu/rpc_fprintf.cpp
+++ /dev/null
@@ -1,75 +0,0 @@
-//===-- GPU implementation of fprintf -------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include "rpc_fprintf.h"
-
-#include "src/__support/CPP/string_view.h"
-#include "src/__support/GPU/utils.h"
-#include "src/__support/RPC/rpc_client.h"
-#include "src/__support/common.h"
-#include "src/__support/macros/config.h"
-#include "src/stdio/gpu/file.h"
-
-namespace LIBC_NAMESPACE_DECL {
-
-template <uint16_t opcode>
-int fprintf_impl(::FILE *__restrict file, const char *__restrict format,
- size_t format_size, void *args, size_t args_size) {
- uint64_t mask = gpu::get_lane_mask();
- rpc::Client::Port port = rpc::client.open<opcode>();
-
- if constexpr (opcode == RPC_PRINTF_TO_STREAM) {
- port.send([&](rpc::Buffer *buffer) {
- buffer->data[0] = reinterpret_cast<uintptr_t>(file);
- });
- }
-
- port.send_n(format, format_size);
- port.recv([&](rpc::Buffer *buffer) {
- args_size = static_cast<size_t>(buffer->data[0]);
- });
- port.send_n(args, args_size);
-
- uint32_t ret = 0;
- for (;;) {
- const char *str = nullptr;
- port.recv([&](rpc::Buffer *buffer) {
- ret = static_cast<uint32_t>(buffer->data[0]);
- str = reinterpret_cast<const char *>(buffer->data[1]);
- });
- // If any lanes have a string argument it needs to be copied back.
- if (!gpu::ballot(mask, str))
- break;
-
- uint64_t size = str ? internal::string_length(str) + 1 : 0;
- port.send_n(str, size);
- }
-
- port.close();
- return ret;
-}
-
-// TODO: Delete this and port OpenMP to use `printf`.
-// place of varargs. Once varargs support is added we will use that to
-// implement the real version.
-LLVM_LIBC_FUNCTION(int, rpc_fprintf,
- (::FILE *__restrict stream, const char *__restrict format,
- void *args, size_t size)) {
- cpp::string_view str(format);
- if (stream == stdout)
- return fprintf_impl<RPC_PRINTF_TO_STDOUT>(stream, format, str.size() + 1,
- args, size);
- else if (stream == stderr)
- return fprintf_impl<RPC_PRINTF_TO_STDERR>(stream, format, str.size() + 1,
- args, size);
- else
- return fprintf_impl<RPC_PRINTF_TO_STREAM>(stream, format, str.size() + 1,
- args, size);
-}
-
-} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/gpu/rpc_fprintf.h b/libc/src/gpu/rpc_fprintf.h
deleted file mode 100644
index 7658b214c07c2..0000000000000
--- a/libc/src/gpu/rpc_fprintf.h
+++ /dev/null
@@ -1,23 +0,0 @@
-//===-- Implementation header for RPC functions -----------------*- C++ -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
-#define LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
-
-#include "hdr/types/FILE.h"
-#include "src/__support/macros/config.h"
-#include <stddef.h>
-
-namespace LIBC_NAMESPACE_DECL {
-
-int rpc_fprintf(::FILE *__restrict stream, const char *__restrict format,
- void *argc, size_t size);
-
-} // namespace LIBC_NAMESPACE_DECL
-
-#endif // LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
index 42a6bac4fa6f2..02b0d436451a3 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
@@ -437,7 +437,8 @@ bool AMDGPUPrintfRuntimeBindingImpl::run(Module &M) {
return false;
auto PrintfFunction = M.getFunction("printf");
- if (!PrintfFunction || !PrintfFunction->isDeclaration())
+ if (!PrintfFunction || !PrintfFunction->isDeclaration() ||
+ M.getModuleFlag("openmp"))
return false;
for (auto &U : PrintfFunction->uses()) {
diff --git a/offload/DeviceRTL/include/LibC.h b/offload/DeviceRTL/include/LibC.h
index dde86af783af9..59a795cc62e0e 100644
--- a/offload/DeviceRTL/include/LibC.h
+++ b/offload/DeviceRTL/include/LibC.h
@@ -18,7 +18,6 @@ extern "C" {
int memcmp(const void *lhs, const void *rhs, size_t count);
void memset(void *dst, int C, size_t count);
-
int printf(const char *format, ...);
}
diff --git a/offload/DeviceRTL/src/LibC.cpp b/offload/DeviceRTL/src/LibC.cpp
index 4bca5d29643fe..291ceb023a69c 100644
--- a/offload/DeviceRTL/src/LibC.cpp
+++ b/offload/DeviceRTL/src/LibC.cpp
@@ -11,44 +11,33 @@
#pragma omp begin declare target device_type(nohost)
namespace impl {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t);
+int32_t omp_vprintf(const char *Format, __builtin_va_list vlist);
}
+#ifndef OMPTARGET_HAS_LIBC
+namespace impl {
#pragma omp begin declare variant match( \
device = {arch(nvptx, nvptx64)}, \
implementation = {extension(match_any)})
-extern "C" int32_t vprintf(const char *, void *);
-namespace impl {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
- return vprintf(Format, Arguments);
+extern "C" int vprintf(const char *format, ...);
+int omp_vprintf(const char *Format, __builtin_va_list vlist) {
+ return vprintf(Format, vlist);
}
-} // namespace impl
#pragma omp end declare variant
#pragma omp begin declare variant match(device = {arch(amdgcn)})
-
-#ifdef OMPTARGET_HAS_LIBC
-// TODO: Remove this handling once we have varargs support.
-extern "C" struct FILE *stdout;
-extern "C" int32_t rpc_fprintf(FILE *, const char *, void *, uint64_t);
-
-namespace impl {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t Size) {
- return rpc_fprintf(stdout, Format, Arguments, Size);
-}
+int omp_vprintf(const char *Format, __builtin_va_list) { return -1; }
+#pragma omp end declare variant
} // namespace impl
-#else
-// We do not have a vprintf implementation for AMD GPU so we use a stub.
-namespace impl {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
- return -1;
+
+extern "C" int printf(const char *Format, ...) {
+ __builtin_va_list vlist;
+ __builtin_va_start(vlist, Format);
+ return impl::omp_vprintf(Format, vlist);
}
-} // namespace impl
-#endif
-#pragma omp end declare variant
+#endif // OMPTARGET_HAS_LIBC
extern "C" {
-
[[gnu::weak]] int memcmp(const void *lhs, const void *rhs, size_t count) {
auto *L = reinterpret_cast<const unsigned char *>(lhs);
auto *R = reinterpret_cast<const unsigned char *>(rhs);
@@ -65,11 +54,6 @@ extern "C" {
for (size_t I = 0; I < count; ++I)
dstc[I] = C;
}
-
-/// printf() calls are rewritten by CGGPUBuiltin to __llvm_omp_vprintf
-int32_t __llvm_omp_vprintf(const char *Format, void *Arguments, uint32_t Size) {
- return impl::omp_vprintf(Format, Arguments, Size);
-}
}
#pragma omp end declare target
More information about the libc-commits
mailing list