diff --git a/openmp/libomptarget/DeviceRTL/CMakeLists.txt b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
--- a/openmp/libomptarget/DeviceRTL/CMakeLists.txt
+++ b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
@@ -85,7 +85,6 @@
${include_directory}/Configuration.h
${include_directory}/Debug.h
${include_directory}/Interface.h
- ${include_directory}/LibC.h
${include_directory}/Mapping.h
${include_directory}/State.h
${include_directory}/Synchronization.h
@@ -97,7 +96,6 @@
${source_directory}/Configuration.cpp
${source_directory}/Debug.cpp
${source_directory}/Kernel.cpp
- ${source_directory}/LibC.cpp
${source_directory}/Mapping.cpp
${source_directory}/Misc.cpp
${source_directory}/Parallelism.cpp
diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -13,12 +13,12 @@
#define OMPTARGET_DEVICERTL_DEBUG_H
#include "Configuration.h"
-#include "LibC.h"
/// Assertion
///
/// {
extern "C" {
+int printf(const char *format, ...);
void __assert_assume(bool condition);
void __assert_fail(const char *assertion, const char *file, unsigned line,
const char *function);
diff --git a/openmp/libomptarget/DeviceRTL/include/LibC.h b/openmp/libomptarget/DeviceRTL/include/LibC.h
deleted file mode 100644
--- a/openmp/libomptarget/DeviceRTL/include/LibC.h
+++ /dev/null
@@ -1,24 +0,0 @@
-//===--------- LibC.h - Simple implementation of libc 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 OMPTARGET_LIBC_H
-#define OMPTARGET_LIBC_H
-
-#include "Types.h"
-
-extern "C" {
-
-int memcmp(const void *lhs, const void *rhs, size_t count);
-
-int printf(const char *format, ...);
-}
-
-#endif
diff --git a/openmp/libomptarget/DeviceRTL/src/LibC.cpp b/openmp/libomptarget/DeviceRTL/src/LibC.cpp
deleted file mode 100644
--- a/openmp/libomptarget/DeviceRTL/src/LibC.cpp
+++ /dev/null
@@ -1,56 +0,0 @@
-//===------- LibC.cpp - Simple implementation of libc 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
-//
-//===----------------------------------------------------------------------===//
-
-#include "LibC.h"
-
-#pragma omp begin declare target device_type(nohost)
-
-namespace impl {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t);
-}
-
-#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);
-}
-} // 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 {
-int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
- return -1;
-}
-} // namespace impl
-#pragma omp end declare variant
-
-extern "C" {
-
-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);
-
- for (size_t I = 0; I < count; ++I)
- if (L[I] != R[I])
- return (int)L[I] - (int)R[I];
-
- return 0;
-}
-
-/// 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
diff --git a/openmp/libomptarget/DeviceRTL/src/Misc.cpp b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
--- a/openmp/libomptarget/DeviceRTL/src/Misc.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
@@ -92,5 +92,37 @@
double omp_get_wtime(void) { return ompx::impl::getWTime(); }
}
+namespace impl {
+int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t);
+}
+
+#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);
+}
+} // 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 {
+int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
+ return -1;
+}
+} // namespace impl
+#pragma omp end declare variant
+
+extern "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
diff --git a/openmp/libomptarget/test/libc/equal.cpp b/openmp/libomptarget/test/libc/equal.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/libc/equal.cpp
@@ -0,0 +1,25 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// REQUIRES: libc
+
+#include <algorithm>
+
+extern "C" int printf(const char *, ...);
+
+// std::equal is lowered to the libc function memcmp.
+void test_memcpy() {
+ int r = 0;
+#pragma omp target map(from : r)
+ {
+ int x[2] = {0, 0};
+ int y[2] = {0, 0};
+ int z[2] = {0, 1};
+ bool eq1 = std::equal(x, x + 2, y);
+ bool eq2 = std::equal(x, x + 2, z);
+ r = eq1 && !eq2;
+ }
+ // CHECK: memcmp: PASS
+ printf("memcmp: %s\n", r ? "PASS" : "FAIL");
+}
+
+int main(int argc, char *argv[]) { test_memcpy(); }