[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

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

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

80

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

 /*===- __clang_openmp_device_functions.h - OpenMP device function declares -===
  *
  * 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_DEVICE_FUNCTIONS_H__
 #define __CLANG_OPENMP_DEVICE_FUNCTIONS_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
+#ifdef __cplusplus
+// Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the
+// need to `include <new>`.
+#include <new>
+#endif
+
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 #define __CUDA__

 #undef __CUDA__
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #pragma omp end declare variant
 
 #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
 
 #include_next <new>
 
-#if defined(__NVPTX__) && defined(_OPENMP)
+#if defined(_OPENMP)
 
 #include <cstdlib>
 
 #pragma push_macro("OPENMP_NOEXCEPT")
 #if __cplusplus >= 201103L
 #define OPENMP_NOEXCEPT noexcept
 #else
 #define OPENMP_NOEXCEPT

clang/test/Headers/Inputs/include/cstdlib

 #pragma once
 
 #include <stdlib.h>
 
+void *malloc(size_t);
+void free(void*);
+
 #if __cplusplus >= 201703L
 extern int abs (int __x) throw()  __attribute__ ((__const__)) ;
 extern long int labs (long int __x) throw() __attribute__ ((__const__)) ;
 extern float fabs (float __x) throw() __attribute__ ((__const__)) ;
 #else
 extern int abs (int __x) __attribute__ ((__const__)) ;
 extern long int labs (long int __x) __attribute__ ((__const__)) ;
 extern float fabs (float __x) __attribute__ ((__const__)) ;

clang/test/Headers/Inputs/include/new

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

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