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 &Args, llvm::SmallVectorImpl &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 +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); +}