[clang] 5d8d994 - [OpenMP] Make sure classes work on the device as they do on the host

[Johannes Doerfert via cfe-commits]

Author: Johannes Doerfert
Date: 2021-05-06T02:10:30-05:00
New Revision: 5d8d994dfbe38fe86b1d883daa9fd8e47cdc1376

URL: https://github.com/llvm/llvm-project/commit/5d8d994dfbe38fe86b1d883daa9fd8e47cdc1376
DIFF: https://github.com/llvm/llvm-project/commit/5d8d994dfbe38fe86b1d883daa9fd8e47cdc1376.diff

LOG: [OpenMP] Make sure classes work on the device as they do on the host

We do provide `operator delete(void*)` in `<new>` but it should be
available by default. This is mostly boilerplate to test it and the
unconditional include of `<new>` in the header we always in include
on the device.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D100620

Added: 
    clang/test/Headers/Inputs/include/new
    clang/test/Headers/target_include_new.cpp

Modified: 
    clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
    clang/lib/Headers/openmp_wrappers/new
    clang/test/Headers/Inputs/include/stdlib.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
index 406c9748e286e..953857badfc4c 100644
--- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
+++ b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
@@ -39,4 +39,46 @@ extern "C" {
 
 #pragma omp end declare variant
 
+// Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the
+// need to `include <new>` in C++ mode.
+#ifdef __cplusplus
+
+// We require malloc/free.
+#include <cstdlib>
+
+#pragma push_macro("OPENMP_NOEXCEPT")
+#if __cplusplus >= 201103L
+#define OPENMP_NOEXCEPT noexcept
+#else
+#define OPENMP_NOEXCEPT
+#endif
+
+// Device overrides for non-placement new and delete.
+inline void *operator new(__SIZE_TYPE__ size) {
+  if (size == 0)
+    size = 1;
+  return ::malloc(size);
+}
+
+inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); }
+
+inline void operator delete(void *ptr)OPENMP_NOEXCEPT { ::free(ptr); }
+
+inline void operator delete[](void *ptr) OPENMP_NOEXCEPT {
+  ::operator delete(ptr);
+}
+
+// Sized delete, C++14 only.
+#if __cplusplus >= 201402L
+inline void operator delete(void *ptr, __SIZE_TYPE__ size)OPENMP_NOEXCEPT {
+  ::operator delete(ptr);
+}
+inline void operator delete[](void *ptr, __SIZE_TYPE__ size) OPENMP_NOEXCEPT {
+  ::operator delete(ptr);
+}
+#endif
+
+#pragma pop_macro("OPENMP_NOEXCEPT")
+#endif
+
 #endif

diff  --git a/clang/lib/Headers/openmp_wrappers/new b/clang/lib/Headers/openmp_wrappers/new
index 1387d925b126d..985ddc567f494 100644
--- a/clang/lib/Headers/openmp_wrappers/new
+++ b/clang/lib/Headers/openmp_wrappers/new
@@ -9,6 +9,8 @@
 #ifndef __CLANG_OPENMP_WRAPPERS_NEW
 #define __CLANG_OPENMP_WRAPPERS_NEW
 
+// We need the system <new> for the std::nothrow_t. The new/delete operators
+// which do not use nothrow_t are provided without the <new> header.
 #include_next <new>
 
 #if defined(__NVPTX__) && defined(_OPENMP)
@@ -22,48 +24,24 @@
 #define OPENMP_NOEXCEPT
 #endif
 
-// Device overrides for non-placement new and delete.
-inline void *operator new(__SIZE_TYPE__ size) {
-  if (size == 0)
-    size = 1;
-  return ::malloc(size);
-}
 inline void *operator new(__SIZE_TYPE__ size,
                           const std::nothrow_t &) OPENMP_NOEXCEPT {
   return ::operator new(size);
 }
 
-inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); }
 inline void *operator new[](__SIZE_TYPE__ size, const std::nothrow_t &) {
   return ::operator new(size);
 }
 
-inline void operator delete(void *ptr)OPENMP_NOEXCEPT {
-  if (ptr)
-    ::free(ptr);
-}
 inline void operator delete(void *ptr, const std::nothrow_t &)OPENMP_NOEXCEPT {
   ::operator delete(ptr);
 }
 
-inline void operator delete[](void *ptr) OPENMP_NOEXCEPT {
-  ::operator delete(ptr);
-}
 inline void operator delete[](void *ptr,
                               const std::nothrow_t &) OPENMP_NOEXCEPT {
   ::operator delete(ptr);
 }
 
-// Sized delete, C++14 only.
-#if __cplusplus >= 201402L
-inline void operator delete(void *ptr, __SIZE_TYPE__ size)OPENMP_NOEXCEPT {
-  ::operator delete(ptr);
-}
-inline void operator delete[](void *ptr, __SIZE_TYPE__ size) OPENMP_NOEXCEPT {
-  ::operator delete(ptr);
-}
-#endif
-
 #pragma pop_macro("OPENMP_NOEXCEPT")
 #endif
 

diff  --git a/clang/test/Headers/Inputs/include/new b/clang/test/Headers/Inputs/include/new
new file mode 100644
index 0000000000000..8159d5527cc3a
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/new
@@ -0,0 +1,7 @@
+
+namespace std
+{
+
+struct nothrow_t { explicit nothrow_t() = default; };
+
+}

diff  --git a/clang/test/Headers/Inputs/include/stdlib.h b/clang/test/Headers/Inputs/include/stdlib.h
index 516e521df7ec7..47cd80ca84f01 100644
--- a/clang/test/Headers/Inputs/include/stdlib.h
+++ b/clang/test/Headers/Inputs/include/stdlib.h
@@ -1,6 +1,9 @@
 #pragma once
 typedef __SIZE_TYPE__ size_t;
 
+void *malloc(size_t);
+void free(void*);
+
 #ifndef __cplusplus
 extern int abs(int __x) __attribute__((__const__));
 #endif

diff  --git a/clang/test/Headers/target_include_new.cpp b/clang/test/Headers/target_include_new.cpp
new file mode 100644
index 0000000000000..108509e249000
--- /dev/null
+++ b/clang/test/Headers/target_include_new.cpp
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64 -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64 -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+// expected-no-diagnostics
+
+// Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the need to `include <new>`.
+
+// CHECK: define {{.*}}_ZdlPv
+
+#ifndef HEADER
+#define HEADER
+
+class Base {
+  public:
+    virtual ~Base() = default;
+};
+
+class Derived : public Base {
+  public:
+    #pragma omp declare target
+    Derived();
+    #pragma omp end declare target
+};
+
+Derived::Derived() { }
+
+int main(void) {
+  #pragma omp target
+  {
+  }
+  return 0;
+}
+#endif