[Openmp-commits] [openmp] a273d17 - [OpenMP][FIX] Do not add implicit argument to device Ctors and Dtors

[Johannes Doerfert via Openmp-commits]

Author: Johannes Doerfert
Date: 2023-11-01T11:18:11-07:00
New Revision: a273d17d4a0b0d621d111d33302172386d796f98

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

LOG: [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.

Added: 
    openmp/libomptarget/test/offloading/ctor_dtor.cpp

Modified: 
    openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp
    openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp
index 8747502065a5937..106e7a68cd3ae3c 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 bbd6acd19bb01ae..d1294405c04b38f 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 000000000000000..46e9dd46a35631a
--- /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);
+}