From a273d17d4a0b0d621d111d33302172386d796f98 Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes@jdoerfert.de>
Date: Wed, 1 Nov 2023 08:26:38 -0700
Subject: [PATCH] [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.
---
 .../PluginInterface/PluginInterface.cpp       |  3 +++
 .../common/PluginInterface/PluginInterface.h  |  6 +++++
 .../test/offloading/ctor_dtor.cpp             | 24 +++++++++++++++++++
 3 files changed, 33 insertions(+)
 create mode 100644 openmp/libomptarget/test/offloading/ctor_dtor.cpp

diff --git a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp
index 8747502065a59..106e7a68cd3ae 100644
--- a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp
+++ b/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);
diff --git a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h b/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h
index bbd6acd19bb01..d1294405c04b3 100644
--- a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h
+++ b/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!");
diff --git a/openmp/libomptarget/test/offloading/ctor_dtor.cpp b/openmp/libomptarget/test/offloading/ctor_dtor.cpp
new file mode 100644
index 0000000000000..46e9dd46a3563
--- /dev/null
+++ b/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);
+}