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

[jdoerfert]

Summary

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.

[JonChesterfield]

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

20

I think there are legitimate use cases for writing code that doesn't include new. Can we add the forms that are meant to be available without the include here, instead of just pulling in all new?

clang/test/Headers/Inputs/include/cstdlib

5 ↗

(On Diff #337987)

These are in stdlib just above, no? Possibly with noexcept tags

[jdoerfert]

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

20

Hm, what forms would you not like to see in the code by default, and why?

clang/test/Headers/Inputs/include/cstdlib

5 ↗

(On Diff #337987)

They are not, this is the sham layer for testing, but you are so far correct that they should be. I'll move them.

[JonChesterfield]

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

20

I would prefer the default includes exactly the expressions/syntax that C++ provides by default, which may mean writing some inline here though I'm perplexed that we need to do something different to clang compiling c++. I'd have to look up which operators are magic and which are in the header.

At least some reluctance is that we don't have a gpu <new>, so this will be whatever host libstdc++/libc++ is in use, and I'd like to avoid including gpu-unaware code on the gpu wherever possible.

[jdoerfert]

Avoid including <new> by not wrapping <new>

[JonChesterfield]

cstdlib may not exist, probably in the same situations as new not existing, but that can probably be left until the problem arises. I think this looks reasonable, and it'll unblock some applications that don't presently build.

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

74

free(nullptr) is well formed, don't think we need the branch here

[JonChesterfield]

We might want a test that includes <new> along with this header, to check they don't conflict with one another

[jdoerfert]

In D100620#2733707, @JonChesterfield wrote:

We might want a test that includes <new> along with this header, to check they don't conflict with one another

For that I need to replicate the potentially conflicting <new> stuff in the test header mockups. I'll try it later on one system though.

[JonChesterfield]

In D100620#2733708, @jdoerfert wrote:

For that I need to replicate the potentially conflicting <new> stuff in the test header mockups. I'll try it later on one system though.

Yep, may not be able to check in something that pulls <new> out of the host system.

It seems inconvenient that clang doesn't automatically inject the new/delete forms that are intended to be available without including new.

[jdoerfert]

Split the new/delete operators based on std::nothrow_t and provide the ones
that do not use this type by default, the others as part of <new>.

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 #define __CUDA__
 #define __OPENMP_NVPTX__
 
 /// Include declarations for libdevice functions.
 #include <__clang_cuda_libdevice_declares.h>
 
 /// Provide definitions for these functions.
 #include <__clang_cuda_device_functions.h>
 
 #undef __OPENMP_NVPTX__
 #undef __CUDA__
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #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

clang/lib/Headers/openmp_wrappers/new

 //===--------- new - OPENMP wrapper for <new> ------------------------------===
 //
 // 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 __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)
 
 #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,
                           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
 
 #endif // include guard

clang/test/Headers/Inputs/include/new

+
+namespace std
+{
+
+struct nothrow_t { explicit nothrow_t() = default; };
+
+}

clang/test/Headers/Inputs/include/stdlib.h

 #pragma once
 typedef __SIZE_TYPE__ size_t;
 
+void *malloc(size_t);
+void free(void*);
+
 #ifndef __cplusplus
 extern int abs(int __x) __attribute__((__const__));
 #endif

clang/test/Headers/target_include_new.cpp

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