[OpenMP][FIX] Do not add implicit argument to device Ctors and Dtors

Constructors and destructors on the device do not take any arguments,
also not the implicit dyn_ptr argument other kernels automatically take.

Author: jdoerfert

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp

@@ -511,6 +511,9 @@ void *GenericKernelTy::prepareArgs(
     uint32_t &NumArgs, llvm::SmallVectorImpl<void *> &Args,
     llvm::SmallVectorImpl<void *> &Ptrs,
     KernelLaunchEnvironmentTy *KernelLaunchEnvironment) const {
+  if (isCtorOrDtor())
+    return nullptr;
+
   NumArgs += 1;
 
   Args.resize(NumArgs);

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h

@@ -284,6 +284,12 @@ struct GenericKernelTy {
   /// Get the kernel name.
   const char *getName() const { return Name; }
 
+  /// Return true if this kernel is a constructor or destructor.
+  bool isCtorOrDtor() const {
+    // TODO: This is not a great solution and should be revisited.
+    return StringRef(Name).endswith("tor");
+  }
+
   /// Get the kernel image.
   DeviceImageTy &getImage() const {
     assert(ImagePtr && "Kernel is not initialized!");

openmp/libomptarget/test/offloading/ctor_dtor.cpp

@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
+//
+#include <cstdio>
+struct S {
+  S() : i(7) {}
+  ~S() { foo(); }
+  int foo() { return i; }
+
+private:
+  int i;
+};
+
+S s;
+#pragma omp declare target(s)
+
+int main() {
+  int r;
+#pragma omp target map(from : r)
+  r = s.foo();
+
+  // CHECK: 7
+  printf("%i\n", r);
+}