[OpenMP][DeviceRTL] Implement libc function `memcmp`

[tianshilei1992]

Summary

We will add some simple implementation of libc functions starting from
this patch, and the first one is memcmp, which is reported in #56929. Note that
malloc and free are not included in this patch because of the use of
declare variant. In the near future we will implement the two functions w/o
using any vendor provided function.

This fixes #56929.

[tianshilei1992]

fix headers

[tianshilei1992]

remove one extra check

[tianshilei1992]

add printf

[tschuett]

https://discourse.llvm.org/t/llvm-embedded-toolchains-working-group-sync-up/63270/7
There was a discussion to use the in-tree libc for GPUs at their last meeting.

[jdoerfert]

In D131182#3700422, @tschuett wrote:

https://discourse.llvm.org/t/llvm-embedded-toolchains-working-group-sync-up/63270/7
There was a discussion to use the in-tree libc for GPUs at their last meeting.

yeah, this will move again soon. For now we need to fix a bug (reported for OpenMP but existing for all GPU + C++ languages).

[tianshilei1992]

In D131182#3700422, @tschuett wrote:

https://discourse.llvm.org/t/llvm-embedded-toolchains-working-group-sync-up/63270/7
There was a discussion to use the in-tree libc for GPUs at their last meeting.

Thanks for the info. Will join in the meeting next time.

[tianshilei1992]

Will push a change to fix the printf issue on AMDGPU.

[ronlieb]

In D131182#3700502, @tianshilei1992 wrote:

Will push a change to fix the printf issue on AMDGPU.

thanks for being on top of this

[tianshilei1992]

In D131182#3700507, @ronlieb wrote:

In D131182#3700502, @tianshilei1992 wrote:

Will push a change to fix the printf issue on AMDGPU.

thanks for being on top of this

https://github.com/llvm/llvm-project/commit/459e3c51840a7f09536a7839deb7184b53bd7e92 should do it.

openmp/libomptarget/DeviceRTL/CMakeLists.txt

 # Activate RTL message dumps if requested by the user.
 set(LIBOMPTARGET_DEVICE_DEBUG FALSE CACHE BOOL
   "Activate DeviceRTL debug messages.")
 
 set(include_files
   ${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
   ${include_directory}/Types.h
   ${include_directory}/Utils.h
 )
 
 set(src_files
   ${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
   ${source_directory}/Reduction.cpp
   ${source_directory}/State.cpp
   ${source_directory}/Synchronization.cpp
   ${source_directory}/Tasking.cpp
   ${source_directory}/Utils.cpp

openmp/libomptarget/DeviceRTL/include/Debug.h

 //===-------- Debug.h ---- Debug utilities ------------------------ 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_DEVICERTL_DEBUG_H
 #define OMPTARGET_DEVICERTL_DEBUG_H
 
 #include "Configuration.h"
+#include "LibC.h"
 
 /// Assertion
 ///
 /// {
 extern "C" {
 void __assert_assume(bool condition);
 void __assert_fail(const char *assertion, const char *file, unsigned line,
                    const char *function);
 }
 
 #define ASSERT(expr)                                                           \
   {                                                                            \
     if (config::isDebugMode(config::DebugKind::Assertion) && !(expr))          \
       __assert_fail(#expr, __FILE__, __LINE__, __PRETTY_FUNCTION__);           \
     else                                                                       \
       __assert_assume(expr);                                                   \
   }
 
 ///}
 
-/// Print
-/// printf() calls are rewritten by CGGPUBuiltin to __llvm_omp_vprintf
-/// {
-
-extern "C" {
-int printf(const char *format, ...);
-}
-
 #define PRINTF(fmt, ...) (void)printf(fmt, ##__VA_ARGS__);
 #define PRINT(str) PRINTF("%s", str)
 
 ///}
 
 /// Enter a debugging scope for performing function traces. Enabled with
 /// FunctionTracting set in the debug kind.
 #define FunctionTracingRAII()                                                  \

openmp/libomptarget/DeviceRTL/include/LibC.h

+//===--------- 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

openmp/libomptarget/DeviceRTL/include/Types.h

 using int8_t = char;
 using uint8_t = unsigned char;
 using int16_t = short;
 using uint16_t = unsigned short;
 using int32_t = int;
 using uint32_t = unsigned int;
 using int64_t = long;
 using uint64_t = unsigned long;
+using size_t = decltype(sizeof(char));
 
 static_assert(sizeof(int8_t) == 1, "type size mismatch");
 static_assert(sizeof(uint8_t) == 1, "type size mismatch");
 static_assert(sizeof(int16_t) == 2, "type size mismatch");
 static_assert(sizeof(uint16_t) == 2, "type size mismatch");
 static_assert(sizeof(int32_t) == 4, "type size mismatch");
 static_assert(sizeof(uint32_t) == 4, "type size mismatch");
 static_assert(sizeof(int64_t) == 8, "type size mismatch");

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();
 }
-
-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)})
-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
-
-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.
 __attribute__((loader_uninitialized)) static uint32_t Level;
 #pragma omp allocate(Level) allocator(omp_pteam_mem_alloc)
 
 DebugEntryRAII::DebugEntryRAII(const char *File, const unsigned Line,
                                const char *Function) {

openmp/libomptarget/DeviceRTL/src/LibC.cpp

+//===------- LibC.c - 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

openmp/libomptarget/DeviceRTL/src/exports

 omp_*
 *llvm_*
 __kmpc_*
+
+memcmp
+printf

openmp/libomptarget/test/offloading/test_libc.cpp

+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <algorithm>
+
+extern "C" int printf(const char *, ...);
+
+// std::equal is lowered to libc function memcmp.
+void test_memcpy() {
+#pragma omp target
+  {
+    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);
+    bool r = eq1 && !eq2;
+    printf("memcmp: %s\n", r ? "PASS" : "FAIL");
+  }
+}
+
+int main(int argc, char *argv[]) {
+  test_memcpy();
+
+  return 0;
+}
+
+// CHECK: memcmp: PASS