[clang] [llvm] [OpenMP][clang] Indirect and Virtual function call mapping from host to device (PR #159857)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 19 15:01:49 PDT 2025
https://github.com/Jason-VanBeusekom created https://github.com/llvm/llvm-project/pull/159857
This adds a feature to insert CPU-to-GPU function pointer translation at GPU
call-sites
How it works:
Search Clang AST within "omp target" region for CallExprs.
Store them in the set CrayOMPTargetCalls (kept by CodeGenModule).
When generating code for function calls that use function pointers,
check whether the associated CallExpr is in the set CrayOMPTargetCalls.
If YES, insert a call to devicertl function __llvm_omp_indirect_call_lookup.
This is the second out of Two PR's to implement this, commit 5247c1f2ad7e1ae05cc92daca7979d7c9d838cb2 is not a part of the Pull request and is handled in: https://github.com/llvm/llvm-project/pull/159856
>From 77fd376b5b87eab76bda14e7e457ea80e8e09f20 Mon Sep 17 00:00:00 2001
From: "jason.van-beusekom at hpe.com" <jason.van-beusekom at hpe.com>
Date: Fri, 12 Sep 2025 14:07:54 -0500
Subject: [PATCH 1/2] [OpenMP][clang] Register Vtables on device for indirect
calls Runtime / Registration support for indirect and virtual function calls
in OpenMP target regions - Register Vtable's to OpenMP offload table - Modify
PluginInterface to register Vtables to indirect call table This Patch does
not have the logic for calling __llvm_omp_indirect_call_lookup, and lacks
implementation logic ---------
Co-authored-by: Chi-Chun Chen <chichunchen844 at gmail.com>
Co-authored-by: Jeffery Sandoval <jeffrey.sandoval at hpe.com>
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 128 ++++++++
clang/lib/CodeGen/CGOpenMPRuntime.h | 13 +
clang/lib/CodeGen/CGStmtOpenMP.cpp | 4 +
clang/lib/CodeGen/CGVTables.cpp | 6 +
clang/lib/CodeGen/CGVTables.h | 4 +
clang/lib/CodeGen/CodeGenModule.h | 3 +
clang/test/OpenMP/target_vtable_codegen.cpp | 280 ++++++++++++++++++
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 5 +-
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 19 +-
offload/include/omptarget.h | 2 +
offload/libomptarget/PluginManager.cpp | 7 +-
offload/libomptarget/device.cpp | 37 ++-
12 files changed, 497 insertions(+), 11 deletions(-)
create mode 100644 clang/test/OpenMP/target_vtable_codegen.cpp
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a503aaf613e30..028d14e897667 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1771,12 +1771,126 @@ void CGOpenMPRuntime::emitDeclareTargetFunction(const FunctionDecl *FD,
Addr->setVisibility(llvm::GlobalValue::ProtectedVisibility);
}
+ // Register the indirect Vtable:
+ // This is similar to OMPTargetGlobalVarEntryIndirect, except that the
+ // size field refers to the size of memory pointed to, not the size of
+ // the pointer symbol itself (which is implicitly the size of a pointer).
OMPBuilder.OffloadInfoManager.registerDeviceGlobalVarEntryInfo(
Name, Addr, CGM.GetTargetTypeStoreSize(CGM.VoidPtrTy).getQuantity(),
llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect,
llvm::GlobalValue::WeakODRLinkage);
}
+void CGOpenMPRuntime::registerVTableOffloadEntry(llvm::GlobalVariable *VTable,
+ const VarDecl *VD) {
+ // TODO: add logic to avoid duplicate vtable registrations per
+ // translation unit; though for external linkage, this should no
+ // longer be an issue - or at least we can avoid the issue by
+ // checking for an existing offloading entry. But, perhaps the
+ // better approach is to defer emission of the vtables and offload
+ // entries until later (by tracking a list of items that need to be
+ // emitted).
+
+ llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+
+ // Generate a new externally visible global to point to the
+ // internally visible vtable. Doing this allows us to keep the
+ // visibility and linkage of the associated vtable unchanged while
+ // allowing the runtime to access its value. The externally
+ // visible global var needs to be emitted with a unique mangled
+ // name that won't conflict with similarly named (internal)
+ // vtables in other translation units.
+
+ // Register vtable with source location of dynamic object in map
+ // clause.
+ llvm::TargetRegionEntryInfo EntryInfo = getEntryInfoFromPresumedLoc(
+ CGM, OMPBuilder, VD->getCanonicalDecl()->getBeginLoc(),
+ VTable->getName());
+
+ llvm::GlobalVariable *Addr = VTable;
+ size_t PointerSize = CGM.getDataLayout().getPointerSize();
+ SmallString<128> AddrName;
+ OMPBuilder.OffloadInfoManager.getTargetRegionEntryFnName(AddrName, EntryInfo);
+ AddrName.append("addr");
+
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
+ Addr = new llvm::GlobalVariable(
+ CGM.getModule(), VTable->getType(),
+ /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, VTable,
+ AddrName,
+ /*InsertBefore*/ nullptr, llvm::GlobalValue::NotThreadLocal,
+ CGM.getModule().getDataLayout().getDefaultGlobalsAddressSpace());
+ Addr->setVisibility(llvm::GlobalValue::ProtectedVisibility);
+ }
+ OMPBuilder.OffloadInfoManager.registerDeviceGlobalVarEntryInfo(
+ AddrName, VTable,
+ CGM.getDataLayout().getTypeAllocSize(VTable->getInitializer()->getType()),
+ llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable,
+ llvm::GlobalValue::WeakODRLinkage);
+}
+
+// Register VTable by scanning through the map clause of OpenMP target region.
+void CGOpenMPRuntime::registerVTable(const OMPExecutableDirective &D) {
+ // Get CXXRecordDecl and VarDecl from Expr.
+ auto getVTableDecl = [](const Expr *E) {
+ QualType VDTy = E->getType();
+ CXXRecordDecl *CXXRecord = nullptr;
+ if (const auto *RefType = VDTy->getAs<LValueReferenceType>())
+ VDTy = RefType->getPointeeType();
+ if (VDTy->isPointerType())
+ CXXRecord = VDTy->getPointeeType()->getAsCXXRecordDecl();
+ else
+ CXXRecord = VDTy->getAsCXXRecordDecl();
+
+ const VarDecl *VD = nullptr;
+ if (auto *DRE = dyn_cast<DeclRefExpr>(E))
+ VD = cast<VarDecl>(DRE->getDecl());
+ return std::pair<CXXRecordDecl *, const VarDecl *>(CXXRecord, VD);
+ };
+
+ // Emit VTable and register the VTable to OpenMP offload entry recursively.
+ std::function<void(CodeGenModule &, CXXRecordDecl *, const VarDecl *)>
+ emitAndRegisterVTable = [&emitAndRegisterVTable](CodeGenModule &CGM,
+ CXXRecordDecl *CXXRecord,
+ const VarDecl *VD) {
+ // Register C++ VTable to OpenMP Offload Entry if it's a new
+ // CXXRecordDecl.
+ if (CXXRecord && CXXRecord->isDynamicClass() &&
+ CGM.getOpenMPRuntime().VTableDeclMap.find(CXXRecord) ==
+ CGM.getOpenMPRuntime().VTableDeclMap.end()) {
+ CGM.getOpenMPRuntime().VTableDeclMap.try_emplace(CXXRecord, VD);
+ CGM.EmitVTable(CXXRecord);
+ auto VTables = CGM.getVTables();
+ auto *VTablesAddr = VTables.GetAddrOfVTable(CXXRecord);
+ if (VTablesAddr) {
+ CGM.getOpenMPRuntime().registerVTableOffloadEntry(VTablesAddr, VD);
+ }
+ // Emit VTable for all the fields containing dynamic CXXRecord
+ for (const FieldDecl *Field : CXXRecord->fields()) {
+ if (CXXRecordDecl *RecordDecl =
+ Field->getType()->getAsCXXRecordDecl()) {
+ emitAndRegisterVTable(CGM, RecordDecl, VD);
+ }
+ }
+ // Emit VTable for all dynamic parent class
+ for (CXXBaseSpecifier &Base : CXXRecord->bases()) {
+ if (CXXRecordDecl *BaseDecl =
+ Base.getType()->getAsCXXRecordDecl()) {
+ emitAndRegisterVTable(CGM, BaseDecl, VD);
+ }
+ }
+ }
+ };
+
+ // Collect VTable from OpenMP map clause.
+ for (const auto *C : D.getClausesOfKind<OMPMapClause>()) {
+ for (const auto *E : C->varlist()) {
+ auto DeclPair = getVTableDecl(E);
+ emitAndRegisterVTable(CGM, DeclPair.first, DeclPair.second);
+ }
+ }
+}
+
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
QualType VarType,
StringRef Name) {
@@ -6249,6 +6363,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr);
}
}
+ registerVTable(D);
}
/// Checks if the expression is constant or does not have non-trivial function
@@ -9955,6 +10070,19 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
if (!S)
return;
+ // Register vtable from device for target data and target directives.
+ // Add this block here since scanForTargetRegionsFunctions ignores
+ // target data by checking if S is a executable directive (target).
+ if (isa<OMPExecutableDirective>(S) &&
+ isOpenMPTargetDataManagementDirective(
+ cast<OMPExecutableDirective>(S)->getDirectiveKind())) {
+ auto &E = *cast<OMPExecutableDirective>(S);
+ // Don't need to check if it's device compile
+ // since scanForTargetRegionsFunctions currently only called
+ // in device compilation.
+ registerVTable(E);
+ }
+
// Codegen OMP target directives that offload compute to the device.
bool RequiresDeviceCodegen =
isa<OMPExecutableDirective>(S) &&
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index eb04eceee236c..0f7937ae95c06 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -605,6 +605,9 @@ class CGOpenMPRuntime {
LValue PosLVal, const OMPTaskDataTy::DependData &Data,
Address DependenciesArray);
+ /// Keep track of VTable Declarations so we don't register duplicate VTable.
+ llvm::DenseMap<CXXRecordDecl*, const VarDecl*> VTableDeclMap;
+
public:
explicit CGOpenMPRuntime(CodeGenModule &CGM);
virtual ~CGOpenMPRuntime() {}
@@ -1111,6 +1114,16 @@ class CGOpenMPRuntime {
virtual void emitDeclareTargetFunction(const FunctionDecl *FD,
llvm::GlobalValue *GV);
+ /// Register VTable to OpenMP offload entry.
+ /// \param VTable VTable of the C++ class.
+ /// \param RD C++ class decl.
+ virtual void registerVTableOffloadEntry(llvm::GlobalVariable *VTable,
+ const VarDecl *VD);
+ /// Emit code for registering vtable by scanning through map clause
+ /// in OpenMP target region.
+ /// \param D OpenMP target directive.
+ virtual void registerVTable(const OMPExecutableDirective &D);
+
/// Creates artificial threadprivate variable with name \p Name and type \p
/// VarType.
/// \param VarType Type of the artificial threadprivate variable.
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index d72cd8fbfd608..582dd0f3ade65 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -7617,6 +7617,10 @@ void CodeGenFunction::EmitOMPUseDeviceAddrClause(
// Generate the instructions for '#pragma omp target data' directive.
void CodeGenFunction::EmitOMPTargetDataDirective(
const OMPTargetDataDirective &S) {
+ // Emit vtable only from host for target data directive.
+ if (!CGM.getLangOpts().OpenMPIsTargetDevice) {
+ CGM.getOpenMPRuntime().registerVTable(S);
+ }
CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true,
/*SeparateBeginEndCalls=*/true);
diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index e14e883a55ac5..de4a67db313ea 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -38,6 +38,12 @@ llvm::Constant *CodeGenModule::GetAddrOfThunk(StringRef Name, llvm::Type *FnTy,
/*DontDefer=*/true, /*IsThunk=*/true);
}
+llvm::GlobalVariable *CodeGenVTables::GetAddrOfVTable(const CXXRecordDecl *RD) {
+ llvm::GlobalVariable *VTable =
+ CGM.getCXXABI().getAddrOfVTable(RD, CharUnits());
+ return VTable;
+}
+
static void setThunkProperties(CodeGenModule &CGM, const ThunkInfo &Thunk,
llvm::Function *ThunkFn, bool ForVTable,
GlobalDecl GD) {
diff --git a/clang/lib/CodeGen/CGVTables.h b/clang/lib/CodeGen/CGVTables.h
index 5c45e355fb145..37458eee02e34 100644
--- a/clang/lib/CodeGen/CGVTables.h
+++ b/clang/lib/CodeGen/CGVTables.h
@@ -122,6 +122,10 @@ class CodeGenVTables {
llvm::GlobalVariable::LinkageTypes Linkage,
const CXXRecordDecl *RD);
+ /// GetAddrOfVTable - Get the address of the VTable for the given record
+ /// decl.
+ llvm::GlobalVariable *GetAddrOfVTable(const CXXRecordDecl *RD);
+
/// EmitThunks - Emit the associated thunks for the given global decl.
void EmitThunks(GlobalDecl GD);
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 3971b296b3f80..4ace1abcb5246 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -754,6 +754,9 @@ class CodeGenModule : public CodeGenTypeCache {
// i32 @__isPlatformVersionAtLeast(i32, i32, i32, i32)
llvm::FunctionCallee IsPlatformVersionAtLeastFn = nullptr;
+ // Store indirect CallExprs that are within an omp target region
+ llvm::SmallPtrSet<const CallExpr *, 16> OMPTargetCalls;
+
InstrProfStats &getPGOStats() { return PGOStats; }
llvm::IndexedInstrProfReader *getPGOReader() const { return PGOReader.get(); }
diff --git a/clang/test/OpenMP/target_vtable_codegen.cpp b/clang/test/OpenMP/target_vtable_codegen.cpp
new file mode 100644
index 0000000000000..276cef4eb8801
--- /dev/null
+++ b/clang/test/OpenMP/target_vtable_codegen.cpp
@@ -0,0 +1,280 @@
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s --check-prefix=CK1
+//
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s --check-prefix=CK2
+//
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s --check-prefix=CK3
+//
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s --check-prefix=CK4
+//
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52 -stdlib=libc++
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 -stdlib=libc++ | FileCheck %s --check-prefix=CK5
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+#ifdef CK1
+
+// Make sure both host and device compilation emit vtable for Dervied
+// CK1-DAG: $_ZN7DerivedD1Ev = comdat any
+// CK1-DAG: $_ZN7DerivedD0Ev = comdat any
+// CK1-DAG: $_ZN7Derived5BaseAEi = comdat any
+// CK1-DAG: $_ZN7Derived8DerivedBEv = comdat any
+// CK1-DAG: $_ZN7DerivedD2Ev = comdat any
+// CK1-DAG: $_ZN4BaseD2Ev = comdat any
+// CK1-DAG: $_ZTV7Derived = comdat any
+class Base {
+public:
+ virtual ~Base() = default;
+ virtual void BaseA(int a) { }
+};
+
+// CK1: @_ZTV7Derived = linkonce_odr unnamed_addr constant { [6 x ptr] }
+class Derived : public Base {
+public:
+ ~Derived() override = default;
+ void BaseA(int a) override { x = a; }
+ virtual void DerivedB() { }
+private:
+ int x;
+};
+
+int main() {
+
+ Derived d;
+ Base& c = d;
+ int a = 50;
+ // Should emit vtable for Derived since d is added to map clause
+#pragma omp target data map (to: d, a)
+ {
+ #pragma omp target map(d)
+ {
+ c.BaseA(a);
+ }
+ }
+ return 0;
+}
+
+#endif // CK1
+
+#ifdef CK2
+
+namespace {
+
+// Make sure both host and device compilation emit vtable for Dervied
+// CK2-DAG: @_ZTVN12_GLOBAL__N_17DerivedE
+// CK2-DAG: @_ZN12_GLOBAL__N_17DerivedD1Ev
+// CK2-DAG: @_ZN12_GLOBAL__N_17DerivedD0Ev
+// CK2-DAG: @_ZN12_GLOBAL__N_17Derived5BaseAEi
+// CK2-DAG: @_ZN12_GLOBAL__N_17Derived8DerivedBEv
+class Base {
+public:
+ virtual ~Base() = default;
+ virtual void BaseA(int a) { }
+};
+
+class Derived : public Base {
+public:
+ ~Derived() override = default;
+ void BaseA(int a) override { x = a; }
+ virtual void DerivedB() { }
+private:
+ int x;
+};
+
+};
+
+int main() {
+
+ Derived d;
+ Base& c = d;
+ int a = 50;
+#pragma omp target data map (to: d, a)
+ {
+ #pragma omp target
+ {
+ c.BaseA(a);
+ }
+ }
+ return 0;
+}
+
+#endif // CK2
+
+#ifdef CK3
+
+// CK3-DAG: @_ZTV6Base_1
+// CK3-DAG: @_ZTV7Derived
+// CK3-DAG: @_ZTV6Base_2
+#pragma omp begin declare target
+
+class Base_1 {
+public:
+ virtual void foo() { }
+ virtual void bar() { }
+};
+
+class Base_2 {
+public:
+ virtual void foo() { }
+ virtual void bar() { }
+};
+
+class Derived : public Base_1, public Base_2 {
+public:
+ virtual void foo() override { }
+ virtual void bar() override { }
+};
+
+#pragma omp end declare target
+
+int main() {
+ Base_1 base;
+ Derived derived;
+
+ // Make sure we emit vtable for parent class (Base_1 and Base_2)
+#pragma omp target data map(derived)
+ {
+ Base_1 *p1 = &derived;
+
+#pragma omp target
+ {
+ p1->foo();
+ p1->bar();
+ }
+ }
+ return 0;
+}
+
+#endif // CK3
+
+#ifdef CK4
+
+// CK4-DAG: @_ZTV3Car
+// CK4-DAG: @_ZTV6Engine
+// CK4-DAG: @_ZTV6Wheels
+// CK4-DAG: @_ZTV7Vehicle
+// CK4-DAG: @_ZTV5Brand
+class Engine {
+public:
+ Engine(const char *type) : type(type) {}
+ virtual ~Engine() {}
+
+ virtual void start() const { }
+
+protected:
+ const char *type;
+};
+
+class Wheels {
+public:
+ Wheels(int count) : count(count) {}
+ virtual ~Wheels() {}
+
+ virtual void roll() const { }
+
+protected:
+ int count;
+};
+
+class Vehicle {
+public:
+ Vehicle(int speed) : speed(speed) {}
+ virtual ~Vehicle() {}
+
+ virtual void move() const { }
+
+protected:
+ int speed;
+};
+
+class Brand {
+public:
+ Brand(const char *brandName) : brandName(brandName) {}
+ virtual ~Brand() {}
+
+ void showBrand() const { }
+
+protected:
+ const char *brandName;
+};
+
+class Car : public Vehicle, public Brand {
+public:
+ Car(const char *brand, int speed, const char *engineType, int wheelCount)
+ : Vehicle(speed), Brand(brand), engine(engineType), wheels(wheelCount) {}
+
+ void move() const override { }
+
+ void drive() const {
+ showBrand();
+ engine.start();
+ wheels.roll();
+ move();
+ }
+
+private:
+ Engine engine;
+ Wheels wheels;
+};
+
+int main() {
+ Car myActualCar("Ford", 100, "Hybrid", 4);
+
+ // Make sure we emit VTable for dynamic class as field
+#pragma omp target map(myActualCar)
+ {
+ myActualCar.drive();
+ }
+ return 0;
+}
+
+#endif // CK4
+
+#ifdef CK5
+
+// CK5-DAG: @_ZTV7Derived
+// CK5-DAG: @_ZTV4Base
+template <typename T>
+class Container {
+private:
+T value;
+public:
+Container() : value() {}
+Container(T val) : value(val) {}
+
+T getValue() const { return value; }
+
+void setValue(T val) { value = val; }
+};
+
+class Base {
+public:
+ virtual void foo() {}
+};
+class Derived : public Base {};
+
+class Test {
+public:
+ Container<Derived> v;
+};
+
+int main() {
+ Test test;
+ Derived d;
+ test.v.setValue(d);
+
+// Make sure we emit VTable for type indirectly (template specialized type)
+#pragma omp target map(test)
+ {
+ test.v.getValue().foo();
+ }
+ return 0;
+}
+
+#endif // CK5
+#endif
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index f43ef932e965a..cc0d4c89f9b9f 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -390,6 +390,8 @@ class OffloadEntriesInfoManager {
OMPTargetGlobalVarEntryIndirect = 0x8,
/// Mark the entry as a register requires global.
OMPTargetGlobalRegisterRequires = 0x10,
+ /// Mark the entry as a declare target indirect vtable.
+ OMPTargetGlobalVarEntryIndirectVTable = 0x20,
};
/// Kind of device clause for declare target variables
@@ -2666,7 +2668,8 @@ class OpenMPIRBuilder {
enum EmitMetadataErrorKind {
EMIT_MD_TARGET_REGION_ERROR,
EMIT_MD_DECLARE_TARGET_ERROR,
- EMIT_MD_GLOBAL_VAR_LINK_ERROR
+ EMIT_MD_GLOBAL_VAR_LINK_ERROR,
+ EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR
};
/// Callback function type
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 220eee3cb8b08..a18a4bcb6d62e 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10246,6 +10246,13 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
continue;
}
break;
+ case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect:
+ case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable:
+ if (!CE->getAddress()) {
+ ErrorFn(EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR, E.second);
+ continue;
+ }
+ break;
default:
break;
}
@@ -10255,12 +10262,17 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
// entry. Indirect variables are handled separately on the device.
if (auto *GV = dyn_cast<GlobalValue>(CE->getAddress()))
if ((GV->hasLocalLinkage() || GV->hasHiddenVisibility()) &&
- Flags != OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect)
+ (Flags !=
+ OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect ||
+ Flags != OffloadEntriesInfoManager::
+ OMPTargetGlobalVarEntryIndirectVTable))
continue;
// Indirect globals need to use a special name that doesn't match the name
// of the associated host global.
- if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect)
+ if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect ||
+ Flags ==
+ OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable)
createOffloadEntry(CE->getAddress(), CE->getAddress(), CE->getVarSize(),
Flags, CE->getLinkage(), CE->getVarName());
else
@@ -10689,7 +10701,8 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
}
return;
}
- if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect)
+ if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect ||
+ Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable)
OffloadEntriesDeviceGlobalVar.try_emplace(VarName, OffloadingEntriesNum,
Addr, VarSize, Flags, Linkage,
VarName.str());
diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h
index 8fd722bb15022..bdcda770f2d37 100644
--- a/offload/include/omptarget.h
+++ b/offload/include/omptarget.h
@@ -94,6 +94,8 @@ enum OpenMPOffloadingDeclareTargetFlags {
OMP_DECLARE_TARGET_INDIRECT = 0x08,
/// This is an entry corresponding to a requirement to be registered.
OMP_REGISTER_REQUIRES = 0x10,
+ /// Mark the entry global as being an indirect vtable.
+ OMP_DECLARE_TARGET_INDIRECT_VTABLE = 0x20,
};
enum TargetAllocTy : int32_t {
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index b57a2f815cba6..0cdeeb2d55f17 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -434,7 +434,8 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
llvm::offloading::EntryTy DeviceEntry = Entry;
if (Entry.Size) {
- if (Device.RTL->get_global(Binary, Entry.Size, Entry.SymbolName,
+ if (!(Entry.Flags & OMP_DECLARE_TARGET_INDIRECT_VTABLE) &&
+ Device.RTL->get_global(Binary, Entry.Size, Entry.SymbolName,
&DeviceEntry.Address) != OFFLOAD_SUCCESS)
REPORT("Failed to load symbol %s\n", Entry.SymbolName);
@@ -443,7 +444,9 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
// the device to point to the memory on the host.
if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
(PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
- if (Device.RTL->data_submit(DeviceId, DeviceEntry.Address,
+ if (!(OMP_DECLARE_TARGET_INDIRECT_VTABLE |
+ OMP_DECLARE_TARGET_INDIRECT) &&
+ Device.RTL->data_submit(DeviceId, DeviceEntry.Address,
Entry.Address,
Entry.Size) != OFFLOAD_SUCCESS)
REPORT("Failed to write symbol for USM %s\n", Entry.SymbolName);
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index 71423ae0c94d9..fa1920eb8e89b 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -112,13 +112,39 @@ setupIndirectCallTable(DeviceTy &Device, __tgt_device_image *Image,
llvm::SmallVector<std::pair<void *, void *>> IndirectCallTable;
for (const auto &Entry : Entries) {
if (Entry.Kind != llvm::object::OffloadKind::OFK_OpenMP ||
- Entry.Size == 0 || !(Entry.Flags & OMP_DECLARE_TARGET_INDIRECT))
+ Entry.Size == 0 ||
+ !(Entry.Flags &
+ (OMP_DECLARE_TARGET_INDIRECT | OMP_DECLARE_TARGET_INDIRECT_VTABLE)))
continue;
- assert(Entry.Size == sizeof(void *) && "Global not a function pointer?");
- auto &[HstPtr, DevPtr] = IndirectCallTable.emplace_back();
-
- void *Ptr;
+ size_t PtrSize = sizeof(void *);
+ if (Entry.Flags & OMP_DECLARE_TARGET_INDIRECT_VTABLE) {
+ // This is a VTable entry, the current entry is the first index of the
+ // VTable and Entry.Size is the total size of the VTable. Unlike the
+ // indirect function case below, the Global is not of size Entry.Size and
+ // is instead of size PtrSize (sizeof(void*)).
+ void *Vtable;
+ void *res;
+ if (Device.RTL->get_global(Binary, PtrSize, Entry.SymbolName, &Vtable))
+ return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
+ "failed to load %s", Entry.SymbolName);
+
+ // HstPtr = Entry.Address;
+ if (Device.retrieveData(&res, Vtable, PtrSize, AsyncInfo))
+ return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
+ "failed to load %s", Entry.SymbolName);
+ // Calculate and emplace entire Vtable from first Vtable byte
+ for (uint64_t i = 0; i < Entry.Size / PtrSize; ++i) {
+ auto &[HstPtr, DevPtr] = IndirectCallTable.emplace_back();
+ HstPtr = (void *)((uintptr_t)Entry.Address + i * PtrSize);
+ DevPtr = (void *)((uintptr_t)res + i * PtrSize);
+ }
+ } else {
+ // Indirect function case: Entry.Size should equal PtrSize since we're
+ // dealing with a single function pointer (not a VTable)
+ assert(Entry.Size == PtrSize && "Global not a function pointer?");
+ auto &[HstPtr, DevPtr] = IndirectCallTable.emplace_back();
+ void *Ptr;
if (Device.RTL->get_global(Binary, Entry.Size, Entry.SymbolName, &Ptr))
return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
"failed to load %s", Entry.SymbolName);
@@ -127,6 +153,7 @@ setupIndirectCallTable(DeviceTy &Device, __tgt_device_image *Image,
if (Device.retrieveData(&DevPtr, Ptr, Entry.Size, AsyncInfo))
return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
"failed to load %s", Entry.SymbolName);
+ }
}
// If we do not have any indirect globals we exit early.
>From 5247c1f2ad7e1ae05cc92daca7979d7c9d838cb2 Mon Sep 17 00:00:00 2001
From: "jason.van-beusekom at hpe.com" <jason.van-beusekom at hpe.com>
Date: Fri, 12 Sep 2025 14:08:24 -0500
Subject: [PATCH 2/2] [OpenMP][clang] Indirect and Virtual function call
mapping from host to device This patch implements the CodeGen logic for
calling __llvm_omp_indirect_call_lookup on the device when an indirect
function call or a virtual function call is made within an OpenMP target
region. --------- Co-authored-by: Youngsuk Kim
---
clang/lib/CodeGen/CGExpr.cpp | 20 +
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 30 ++
clang/lib/CodeGen/ItaniumCXXABI.cpp | 18 +
...target_vtable_omp_indirect_call_lookup.cpp | 51 +++
offload/test/api/omp_indirect_func_basic.c | 97 ++++
offload/test/api/omp_indirect_func_struct.c | 213 +++++++++
offload/test/api/omp_virtual_func.cpp | 161 +++++++
...p_virtual_func_multiple_inheritance_01.cpp | 416 +++++++++++++++++
...p_virtual_func_multiple_inheritance_02.cpp | 428 ++++++++++++++++++
.../test/api/omp_virtual_func_reference.cpp | 80 ++++
10 files changed, 1514 insertions(+)
create mode 100644 clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
create mode 100644 offload/test/api/omp_indirect_func_basic.c
create mode 100644 offload/test/api/omp_indirect_func_struct.c
create mode 100644 offload/test/api/omp_virtual_func.cpp
create mode 100644 offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
create mode 100644 offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
create mode 100644 offload/test/api/omp_virtual_func_reference.cpp
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index e6e4947882544..cc4c21a719f4c 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6583,6 +6583,26 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
Address(Handle, Handle->getType(), CGM.getPointerAlign()));
Callee.setFunctionPointer(Stub);
}
+
+ // Check whether the associated CallExpr is in the set OMPTargetCalls.
+ // If YES, insert a call to devicertl function __llvm_omp_indirect_call_lookup
+ //
+ // This is used for the indriect function Case, virtual function case is
+ // handled in ItaniumCXXABI.cpp
+ if (getLangOpts().OpenMPIsTargetDevice && CGM.OMPTargetCalls.contains(E)) {
+ auto *PtrTy = CGM.VoidPtrTy;
+ llvm::Type *RtlFnArgs[] = {PtrTy};
+ llvm::FunctionCallee DeviceRtlFn = CGM.CreateRuntimeFunction(
+ llvm::FunctionType::get(PtrTy, RtlFnArgs, false),
+ "__llvm_omp_indirect_call_lookup");
+ llvm::Value *Func = Callee.getFunctionPointer();
+ llvm::Type *BackupTy = Func->getType();
+ Func = Builder.CreatePointerBitCastOrAddrSpaceCast(Func, PtrTy);
+ Func = EmitRuntimeCall(DeviceRtlFn, {Func});
+ Func = Builder.CreatePointerBitCastOrAddrSpaceCast(Func, BackupTy);
+ Callee.setFunctionPointer(Func);
+ }
+
llvm::CallBase *LocalCallOrInvoke = nullptr;
RValue Call = EmitCall(FnInfo, Callee, ReturnValue, Args, &LocalCallOrInvoke,
E == MustTailCall, E->getExprLoc());
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 028d14e897667..ac1d467affc00 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -24,6 +24,7 @@
#include "clang/AST/OpenMPClause.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/AST/StmtVisitor.h"
+#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/Basic/OpenMPKinds.h"
#include "clang/Basic/SourceManager.h"
#include "clang/CodeGen/ConstantInitBuilder.h"
@@ -6335,6 +6336,25 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+ class OMPTargetCallCollector
+ : public RecursiveASTVisitor<OMPTargetCallCollector> {
+ public:
+ OMPTargetCallCollector(CodeGenFunction &CGF,
+ llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls)
+ : CGF(CGF), TargetCalls(TargetCalls) {}
+
+ bool VisitCallExpr(CallExpr *CE) {
+ if (!CE->getDirectCallee()) {
+ TargetCalls.insert(CE);
+ }
+ return true;
+ }
+
+ private:
+ CodeGenFunction &CGF;
+ llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls;
+ };
+
llvm::TargetRegionEntryInfo EntryInfo =
getEntryInfoFromPresumedLoc(CGM, OMPBuilder, D.getBeginLoc(), ParentName);
@@ -6343,6 +6363,16 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
[&CGF, &D, &CodeGen](StringRef EntryFnName) {
const CapturedStmt &CS = *D.getCapturedStmt(OMPD_target);
+ // Search Clang AST within "omp target" region for CallExprs.
+ // Store them in the set OMPTargetCalls (kept by CodeGenModule).
+ // This is used for the translation of indirect function calls.
+ const auto &LangOpts = CGF.getLangOpts();
+ if (LangOpts.OpenMPIsTargetDevice) {
+ // Search AST for target "CallExpr"s of "OMPTargetAutoLookup".
+ OMPTargetCallCollector Visitor(CGF, CGF.CGM.OMPTargetCalls);
+ Visitor.TraverseStmt(const_cast<Stmt*>(CS.getCapturedStmt()));
+ }
+
CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
return CGF.GenerateOpenMPCapturedStmtFunction(CS, D);
diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp
index 7dc2eaf1e9f75..1dbfe23cef127 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -2261,6 +2261,24 @@ CGCallee ItaniumCXXABI::getVirtualFunctionPointer(CodeGenFunction &CGF,
llvm::Type *PtrTy = CGM.GlobalsInt8PtrTy;
auto *MethodDecl = cast<CXXMethodDecl>(GD.getDecl());
llvm::Value *VTable = CGF.GetVTablePtr(This, PtrTy, MethodDecl->getParent());
+ /*
+ * For the translate of virtual functions we need to map the (potential) host vtable
+ * to the device vtable. This is done by calling the runtime function
+ * __llvm_omp_indirect_call_lookup.
+ */
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
+ auto *NewPtrTy = CGM.VoidPtrTy;
+ llvm::Type *RtlFnArgs[] = {NewPtrTy};
+ llvm::FunctionCallee DeviceRtlFn = CGM.CreateRuntimeFunction(
+ llvm::FunctionType::get(NewPtrTy, RtlFnArgs, false),
+ "__llvm_omp_indirect_call_lookup");
+ auto *BackupTy = VTable->getType();
+ // Need to convert to generic address space
+ VTable = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(VTable, NewPtrTy);
+ VTable = CGF.EmitRuntimeCall(DeviceRtlFn, {VTable});
+ // convert to original address space
+ VTable = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(VTable, BackupTy);
+ }
uint64_t VTableIndex = CGM.getItaniumVTableContext().getMethodVTableIndex(GD);
llvm::Value *VFunc, *VTableSlotPtr = nullptr;
diff --git a/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp b/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
new file mode 100644
index 0000000000000..52bbb382fb853
--- /dev/null
+++ b/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
@@ -0,0 +1,51 @@
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s --check-prefix=CK1
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+#ifdef CK1
+
+#pragma omp begin declare target
+
+class Base {
+public:
+ virtual int foo() { return 1; }
+ virtual int bar() { return 2; }
+};
+
+class Derived : public Base {
+public:
+ virtual int foo() { return 3; }
+ virtual int bar() { return 4; }
+};
+
+#pragma omp end declare target
+
+int main() {
+ Base base;
+ Derived derived;
+ {
+#pragma omp target data map(base, derived)
+ {
+ Base *pointer1 = &base;
+ Base *pointer2 = &derived;
+
+#pragma omp target
+ {
+ // CK1-DAG: call ptr @__llvm_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__llvm_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__llvm_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__llvm_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ int result1 = pointer1->foo();
+ int result2 = pointer1->bar();
+ int result3 = pointer2->foo();
+ int result4 = pointer2->bar();
+ }
+ }
+ }
+ return 0;
+}
+
+#endif
+#endif
diff --git a/offload/test/api/omp_indirect_func_basic.c b/offload/test/api/omp_indirect_func_basic.c
new file mode 100644
index 0000000000000..ff517247d4932
--- /dev/null
+++ b/offload/test/api/omp_indirect_func_basic.c
@@ -0,0 +1,97 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#define TEST_VAL 5
+
+#pragma omp declare target indirect
+__attribute__((noinline)) __attribute__((optnone)) int direct(int x) {
+ return 2 * x;
+}
+__attribute__((noinline)) __attribute__((optnone)) int indirect_base(int x) {
+ return -1 * x;
+}
+#pragma omp end declare target
+
+int (*indirect)(int) = indirect_base;
+
+void set_indirect_func() { indirect = direct; }
+
+void test_implicit_mapping() {
+ int direct_res, indirect_res;
+
+// Test with initial indirect function pointer (points to indirect_base)
+#pragma omp target map(from : direct_res, indirect_res)
+ {
+ direct_res = direct(TEST_VAL);
+ indirect_res = indirect(TEST_VAL);
+ }
+
+ assert(direct_res == TEST_VAL * 2 &&
+ "Error: direct function returned invalid value");
+ assert(indirect_res == TEST_VAL * -1 &&
+ indirect_res == indirect_base(TEST_VAL) &&
+ "Error: indirect function pointer did not return correct value");
+
+ // Set indirect to point to direct function
+ set_indirect_func();
+
+// Test after setting indirect function pointer
+#pragma omp target map(from : direct_res, indirect_res)
+ {
+ direct_res = direct(TEST_VAL);
+ indirect_res = indirect(TEST_VAL);
+ }
+
+ assert(direct_res == TEST_VAL * 2 &&
+ "Error: direct function returned invalid value");
+ assert(indirect_res == direct_res &&
+ "Error: indirect function pointer did not return correct value after "
+ "being set");
+}
+
+void test_explicit_mapping() {
+ // Reset indirect to initial state
+ indirect = indirect_base;
+
+ int direct_res, indirect_res;
+
+// Test with initial indirect function pointer (points to indirect_base)
+#pragma omp target map(indirect) map(from : direct_res, indirect_res)
+ {
+ direct_res = direct(TEST_VAL);
+ indirect_res = indirect(TEST_VAL);
+ }
+
+ assert(direct_res == TEST_VAL * 2 &&
+ "Error: direct function returned invalid value");
+ assert(indirect_res == TEST_VAL * -1 &&
+ indirect_res == indirect_base(TEST_VAL) &&
+ "Error: indirect function pointer did not return correct value");
+
+ // Set indirect to point to direct function
+ set_indirect_func();
+
+// Test after setting indirect function pointer
+#pragma omp target map(indirect) map(from : direct_res, indirect_res)
+ {
+ direct_res = direct(TEST_VAL);
+ indirect_res = indirect(TEST_VAL);
+ }
+
+ assert(direct_res == TEST_VAL * 2 &&
+ "Error: direct function returned invalid value");
+ assert(indirect_res == direct_res &&
+ "Error: indirect function pointer did not return correct value after "
+ "being set");
+}
+
+int main() {
+ test_implicit_mapping();
+ test_explicit_mapping();
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_indirect_func_struct.c b/offload/test/api/omp_indirect_func_struct.c
new file mode 100644
index 0000000000000..cc2eeb86a2e5c
--- /dev/null
+++ b/offload/test/api/omp_indirect_func_struct.c
@@ -0,0 +1,213 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+#include <assert.h>
+
+#define TEST_VAL 5
+
+#pragma omp declare target indirect
+__attribute__((noinline)) __attribute__((optnone)) int direct_arg(int x) { return 2 * x; }
+__attribute__((noinline)) __attribute__((optnone)) int indirect_base_arg(int x) { return -1 * x; }
+__attribute__((noinline)) __attribute__((optnone)) int direct() { return TEST_VAL; }
+__attribute__((noinline)) __attribute__((optnone)) int indirect_base() { return -1 * TEST_VAL; }
+#pragma omp end declare target
+
+struct indirect_stru {
+ int buffer;
+ int (*indirect1)();
+ int (*indirect0)(int);
+};
+typedef struct {
+ int buffer;
+ int (*indirect1_ptr)();
+ int (*indirect0_ptr)(int);
+} indirect_stru_mapped;
+
+#pragma omp declare mapper (indirect_stru_mapped s) map(s,s.indirect0_ptr,s.indirect1_ptr)
+
+struct indirect_stru global_indirect_val = { .indirect0 = indirect_base_arg, .indirect1 = indirect_base};
+indirect_stru_mapped global_mapped_val = { .indirect0_ptr = indirect_base_arg, .indirect1_ptr = indirect_base};
+
+void test_global_struct_explicit_mapping() {
+ int indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
+ int indirect0_base = indirect_base_arg(TEST_VAL);
+
+ int indirect1_ret = global_indirect_val.indirect1();
+ int indirect1_base = indirect_base();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(global_indirect_val,global_indirect_val.indirect1,global_indirect_val.indirect0) map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
+ indirect1_ret = global_indirect_val.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+
+ global_indirect_val.indirect0 = direct_arg;
+ global_indirect_val.indirect1 = direct;
+
+ indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
+ indirect0_base = direct_arg(TEST_VAL);
+
+ indirect1_ret = global_indirect_val.indirect1();
+ indirect1_base = direct();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(global_indirect_val,global_indirect_val.indirect0,global_indirect_val.indirect1) map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
+ indirect1_ret = global_indirect_val.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+}
+
+void test_local_struct_explicit_mapping() {
+ struct indirect_stru local_indirect_val;
+ local_indirect_val.indirect0 = indirect_base_arg;
+ local_indirect_val.indirect1 = indirect_base;
+
+ int indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
+ int indirect0_base = indirect_base_arg(TEST_VAL);
+
+ int indirect1_ret = local_indirect_val.indirect1();
+ int indirect1_base = indirect_base();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(local_indirect_val,local_indirect_val.indirect1,local_indirect_val.indirect0) map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
+ indirect1_ret = local_indirect_val.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+
+ local_indirect_val.indirect0 = direct_arg;
+ local_indirect_val.indirect1 = direct;
+
+ indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
+ indirect0_base = direct_arg(TEST_VAL);
+
+ indirect1_ret = local_indirect_val.indirect1();
+ indirect1_base = direct();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(local_indirect_val,local_indirect_val.indirect0,local_indirect_val.indirect1) map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
+ indirect1_ret = local_indirect_val.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+}
+
+void test_global_struct_user_mapper() {
+ int indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
+ int indirect0_base = indirect_base_arg(TEST_VAL);
+
+ int indirect1_ret = global_mapped_val.indirect1_ptr();
+ int indirect1_base = indirect_base();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect1_ret = global_mapped_val.indirect1_ptr();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+
+ global_mapped_val.indirect0_ptr = direct_arg;
+ global_mapped_val.indirect1_ptr = direct;
+
+ indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect0_base = direct_arg(TEST_VAL);
+
+ indirect1_ret = global_mapped_val.indirect1_ptr();
+ indirect1_base = direct();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect1_ret = global_mapped_val.indirect1_ptr();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+}
+
+void test_local_struct_user_mapper() {
+ indirect_stru_mapped local_mapped_val;
+ local_mapped_val.indirect0_ptr = indirect_base_arg;
+ local_mapped_val.indirect1_ptr = indirect_base;
+
+ int indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
+ int indirect0_base = indirect_base_arg(TEST_VAL);
+
+ int indirect1_ret = local_mapped_val.indirect1_ptr();
+ int indirect1_base = indirect_base();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect1_ret = local_mapped_val.indirect1_ptr();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+
+ local_mapped_val.indirect0_ptr = direct_arg;
+ local_mapped_val.indirect1_ptr = direct;
+
+ indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect0_base = direct_arg(TEST_VAL);
+
+ indirect1_ret = local_mapped_val.indirect1_ptr();
+ indirect1_base = direct();
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on host");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on host");
+
+ #pragma omp target map(from:indirect0_ret,indirect1_ret)
+ {
+ indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
+ indirect1_ret = local_mapped_val.indirect1_ptr();
+ }
+
+ assert(indirect0_ret == indirect0_base && "Error: indirect0 function pointer returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base && "Error: indirect1 function pointer returned incorrect value on device");
+}
+
+int main() {
+ test_global_struct_explicit_mapping();
+ test_local_struct_explicit_mapping();
+ test_global_struct_user_mapper();
+ test_local_struct_user_mapper();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_virtual_func.cpp b/offload/test/api/omp_virtual_func.cpp
new file mode 100644
index 0000000000000..1cfcb6f4d3a54
--- /dev/null
+++ b/offload/test/api/omp_virtual_func.cpp
@@ -0,0 +1,161 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#define TEST_VAL 10
+
+#pragma omp declare target
+
+class Base {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo() {
+ return 1;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) virtual int bar() {
+ return 2;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo_with_arg(int x) {
+ return x;
+ }
+};
+
+class Derived : public Base {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo() {
+ return 10;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) virtual int bar() {
+ return 20;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo_with_arg(int x) {
+ return -x;
+ }
+};
+
+#pragma omp end declare target
+
+int test_virtual_implicit_map() {
+ Base base;
+ Derived derived;
+ int result1, result2, result3, result4, result5, result6;
+
+ // map both base and derived objects up front, since the spec
+ // requires that when first mapping a C++ object that the static
+ // type must match the dynamic type
+#pragma omp target data map(base, derived)
+ {
+ Base *p1 = &base;
+ Base *p2 = &derived;
+
+#pragma omp target map(from : result1, result2, result3, result4, result5, \
+ result6)
+ {
+ // These calls will fail if Clang does not
+ // translate/attach the vtable pointer in each object
+ result1 = p1->foo();
+ result2 = p1->bar();
+ result3 = p2->foo();
+ result4 = p2->bar();
+ result5 = base.foo();
+ result6 = derived.foo();
+ }
+ }
+
+ assert(result1 == 1 && "p1->foo() implicit map Failed");
+ assert(result2 == 2 && "p1->bar() implicit map Failed");
+ assert(result3 == 10 && "p2->foo() implicit map Failed");
+ assert(result4 == 20 && "p2->bar() implicit map Failed");
+ assert(result5 == 1 && "base.foo() implicit map Failed");
+ assert(result6 == 10 && "derived.foo() implicit map Failed");
+ return 0;
+}
+
+int test_virtual_explicit_map() {
+ Base base;
+ Derived derived;
+ int result1, result2, result3, result4;
+
+ // map both base and derived objects up front, since the spec
+ // requires that when first mapping a C++ object that the static
+ // type must match the dynamic type
+#pragma omp target data map(base, derived)
+ {
+ Base *p1 = &base;
+ Base *p2 = &derived;
+
+#pragma omp target map(p1[0 : 0], p2[0 : 0]) \
+ map(from : result1, result2, result3, result4)
+ {
+ result1 = p1->foo();
+ result2 = p1->bar();
+ result3 = p2->foo();
+ result4 = p2->bar();
+ }
+ }
+
+ assert(result1 == 1 && "p1->foo() explicit map Failed");
+ assert(result2 == 2 && "p1->bar() explicit map Failed");
+ assert(result3 == 10 && "p2->foo() explicit map Failed");
+ assert(result4 == 20 && "p2->bar() explicit map Failed");
+ return 0;
+}
+
+int test_virtual_reference() {
+ Derived ddd;
+ Base cont;
+ Base &bbb = ddd;
+
+ int b_ret, d_ret, c_ret;
+
+#pragma omp target data map(to : ddd, cont)
+ {
+#pragma omp target map(bbb, ddd, cont) map(from : b_ret, d_ret, c_ret)
+ {
+ b_ret = bbb.foo_with_arg(TEST_VAL);
+ d_ret = ddd.foo_with_arg(TEST_VAL);
+ c_ret = cont.foo_with_arg(TEST_VAL);
+ }
+ }
+
+ assert(c_ret == TEST_VAL && "Control Base call failed on gpu");
+ assert(b_ret == -TEST_VAL && "Reference to derived call failed on gpu");
+ assert(d_ret == -TEST_VAL && "Derived call failed on gpu");
+
+ return 0;
+}
+
+int test_virtual_reference_implicit() {
+ Derived ddd;
+ Base cont;
+ Base &bbb = ddd;
+
+ int b_ret, d_ret, c_ret;
+
+#pragma omp target data map(to : ddd, cont)
+ {
+#pragma omp target map(from : b_ret, d_ret, c_ret)
+ {
+ b_ret = bbb.foo_with_arg(TEST_VAL);
+ d_ret = ddd.foo_with_arg(TEST_VAL);
+ c_ret = cont.foo_with_arg(TEST_VAL);
+ }
+ }
+
+ assert(c_ret == TEST_VAL && "Control Base call failed on gpu (implicit)");
+ assert(b_ret == -TEST_VAL && "Reference to derived call failed on gpu (implicit)");
+ assert(d_ret == -TEST_VAL && "Derived call failed on gpu (implicit)");
+
+ return 0;
+}
+
+int main() {
+ test_virtual_implicit_map();
+ test_virtual_explicit_map();
+ test_virtual_reference();
+ test_virtual_reference_implicit();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
new file mode 100644
index 0000000000000..20ab90cd35a3b
--- /dev/null
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
@@ -0,0 +1,416 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp declare target
+
+class Mother {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ MotherFoo(int x) {
+ return x;
+ }
+};
+
+class Father {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ FatherFoo(int x) {
+ return x * 2;
+ }
+};
+
+class Child_1 : public Mother, public Father {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ FatherFoo(int x) {
+ return x * 3;
+ }
+};
+
+class Child_2 : public Mother, public Father {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ MotherFoo(int x) {
+ return x * 4;
+ }
+};
+
+class Child_3 : public Mother, public Father {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ MotherFoo(int x) {
+ return x * 5;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ FatherFoo(int x) {
+ return x * 6;
+ }
+};
+
+#pragma omp end declare target
+
+int test_multiple_inheritance() {
+ Mother mother;
+ Father father;
+ Child_1 child_1;
+ Child_2 child_2;
+ Child_3 child_3;
+
+ // map results back to host
+ int result_mother, result_father;
+ int result_child1_father, result_child1_mother, result_child1_as_mother,
+ result_child1_as_father;
+ int result_child2_mother, result_child2_father, result_child2_as_mother,
+ result_child2_as_father;
+ int result_child3_mother, result_child3_father, result_child3_as_mother,
+ result_child3_as_father;
+
+ // Add reference-based results
+ int ref_result_mother, ref_result_father;
+ int ref_result_child1_father, ref_result_child1_mother,
+ ref_result_child1_as_mother, ref_result_child1_as_father;
+ int ref_result_child2_mother, ref_result_child2_father,
+ ref_result_child2_as_mother, ref_result_child2_as_father;
+ int ref_result_child3_mother, ref_result_child3_father,
+ ref_result_child3_as_mother, ref_result_child3_as_father;
+
+#pragma omp target data map(father, mother, child_1, child_2, child_3)
+ {
+ // Base class pointers and references
+ Mother *ptr_mother = &mother;
+ Father *ptr_father = &father;
+ Mother &ref_mother = mother;
+ Father &ref_father = father;
+
+ // Child_1 pointers, references and casts
+ Child_1 *ptr_child_1 = &child_1;
+ Mother *ptr_child_1_cast_mother = &child_1;
+ Father *ptr_child_1_cast_father = &child_1;
+ Child_1 &ref_child_1 = child_1;
+ Mother &ref_child_1_cast_mother = child_1;
+ Father &ref_child_1_cast_father = child_1;
+
+ // Child_2 pointers, references and casts
+ Child_2 *ptr_child_2 = &child_2;
+ Mother *ptr_child_2_cast_mother = &child_2;
+ Father *ptr_child_2_cast_father = &child_2;
+ Child_2 &ref_child_2 = child_2;
+ Mother &ref_child_2_cast_mother = child_2;
+ Father &ref_child_2_cast_father = child_2;
+
+ // Child_3 pointers and casts
+ Child_3 *ptr_child_3 = &child_3;
+ Mother *ptr_child_3_cast_mother = &child_3;
+ Father *ptr_child_3_cast_father = &child_3;
+ Child_3 &ref_child_3 = child_3;
+ Mother &ref_child_3_cast_mother = child_3;
+ Father &ref_child_3_cast_father = child_3;
+
+#pragma omp target map( \
+ from : result_mother, result_father, result_child1_father, \
+ result_child1_mother, result_child1_as_mother, \
+ result_child1_as_father, result_child2_mother, \
+ result_child2_father, result_child2_as_mother, \
+ result_child2_as_father, result_child3_mother, \
+ result_child3_father, result_child3_as_mother, \
+ result_child3_as_father, ref_result_mother, ref_result_father, \
+ ref_result_child1_father, ref_result_child1_mother, \
+ ref_result_child1_as_mother, ref_result_child1_as_father, \
+ ref_result_child2_mother, ref_result_child2_father, \
+ ref_result_child2_as_mother, ref_result_child2_as_father, \
+ ref_result_child3_mother, ref_result_child3_father, \
+ ref_result_child3_as_mother, ref_result_child3_as_father) \
+ map(ptr_mother[0 : 0], ptr_father[0 : 0], ptr_child_1[0 : 0], \
+ ptr_child_1_cast_mother[0 : 0], ptr_child_1_cast_father[0 : 0], \
+ ptr_child_2[0 : 0], ptr_child_2_cast_mother[0 : 0], \
+ ptr_child_2_cast_father[0 : 0], ptr_child_3[0 : 0], \
+ ptr_child_3_cast_mother[0 : 0], ptr_child_3_cast_father[0 : 0], \
+ ref_mother, ref_father, ref_child_1, ref_child_1_cast_mother, \
+ ref_child_1_cast_father, ref_child_2, ref_child_2_cast_mother, \
+ ref_child_2_cast_father, ref_child_3, ref_child_3_cast_mother, \
+ ref_child_3_cast_father)
+ {
+ // These calls will fail if Clang does not
+ // translate/attach the vtable pointer in each object
+
+ // Pointer-based calls
+ // Mother
+ result_mother = ptr_mother->MotherFoo(1);
+ // Father
+ result_father = ptr_father->FatherFoo(1);
+ // Child_1
+ result_child1_father = ptr_child_1->FatherFoo(1);
+ result_child1_mother = ptr_child_1->MotherFoo(1);
+ result_child1_as_mother = ptr_child_1_cast_mother->MotherFoo(1);
+ result_child1_as_father = ptr_child_1_cast_father->FatherFoo(1);
+ // Child_2
+ result_child2_mother = ptr_child_2->MotherFoo(1);
+ result_child2_father = ptr_child_2->FatherFoo(1);
+ result_child2_as_mother = ptr_child_2_cast_mother->MotherFoo(1);
+ result_child2_as_father = ptr_child_2_cast_father->FatherFoo(1);
+ // Child_3
+ result_child3_mother = ptr_child_3->MotherFoo(1);
+ result_child3_father = ptr_child_3->FatherFoo(1);
+ result_child3_as_mother = ptr_child_3_cast_mother->MotherFoo(1);
+ result_child3_as_father = ptr_child_3_cast_father->FatherFoo(1);
+
+ // Reference-based calls
+ // Mother
+ ref_result_mother = ref_mother.MotherFoo(1);
+ // Father
+ ref_result_father = ref_father.FatherFoo(1);
+ // Child_1
+ ref_result_child1_father = ref_child_1.FatherFoo(1);
+ ref_result_child1_mother = ref_child_1.MotherFoo(1);
+ ref_result_child1_as_mother = ref_child_1_cast_mother.MotherFoo(1);
+ ref_result_child1_as_father = ref_child_1_cast_father.FatherFoo(1);
+ // Child_2
+ ref_result_child2_mother = ref_child_2.MotherFoo(1);
+ ref_result_child2_father = ref_child_2.FatherFoo(1);
+ ref_result_child2_as_mother = ref_child_2_cast_mother.MotherFoo(1);
+ ref_result_child2_as_father = ref_child_2_cast_father.FatherFoo(1);
+ // Child_3
+ ref_result_child3_mother = ref_child_3.MotherFoo(1);
+ ref_result_child3_father = ref_child_3.FatherFoo(1);
+ ref_result_child3_as_mother = ref_child_3_cast_mother.MotherFoo(1);
+ ref_result_child3_as_father = ref_child_3_cast_father.FatherFoo(1);
+ }
+ }
+
+ // Check pointer-based results
+ assert(result_mother == 1 && "Mother Foo failed");
+ assert(result_father == 2 && "Father Foo failed");
+ assert(result_child1_father == 3 && "Child_1 Father Foo failed");
+ assert(result_child1_mother == 1 && "Child_1 Mother Foo failed");
+ assert(result_child1_as_mother == 1 &&
+ "Child_1 Mother Parent Cast Foo failed");
+ assert(result_child1_as_father == 3 &&
+ "Child_1 Father Parent Cast Foo failed");
+ assert(result_child2_mother == 4 && "Child_2 Mother Foo failed");
+ assert(result_child2_father == 2 && "Child_2 Father Foo failed");
+ assert(result_child2_as_mother == 4 &&
+ "Child_2 Mother Parent Cast Foo failed");
+ assert(result_child2_as_father == 2 &&
+ "Child_2 Father Parent Cast Foo failed");
+ assert(result_child3_mother == 5 && "Child_3 Mother Foo failed");
+ assert(result_child3_father == 6 && "Child_3 Father Foo failed");
+ assert(result_child3_as_mother == 5 &&
+ "Child_3 Mother Parent Cast Foo failed");
+ assert(result_child3_as_father == 6 &&
+ "Child_3 Father Parent Cast Foo failed");
+
+ // Check reference-based results
+ assert(ref_result_mother == 1 && "Reference Mother Foo failed");
+ assert(ref_result_father == 2 && "Reference Father Foo failed");
+ assert(ref_result_child1_father == 3 &&
+ "Reference Child_1 Father Foo failed");
+ assert(ref_result_child1_mother == 1 &&
+ "Reference Child_1 Mother Foo failed");
+ assert(ref_result_child1_as_mother == 1 &&
+ "Reference Child_1 Mother Parent Cast Foo failed");
+ assert(ref_result_child1_as_father == 3 &&
+ "Reference Child_1 Father Parent Cast Foo failed");
+ assert(ref_result_child2_mother == 4 &&
+ "Reference Child_2 Mother Foo failed");
+ assert(ref_result_child2_father == 2 &&
+ "Reference Child_2 Father Foo failed");
+ assert(ref_result_child2_as_mother == 4 &&
+ "Reference Child_2 Mother Parent Cast Foo failed");
+ assert(ref_result_child2_as_father == 2 &&
+ "Reference Child_2 Father Parent Cast Foo failed");
+ assert(ref_result_child3_mother == 5 &&
+ "Reference Child_3 Mother Foo failed");
+ assert(ref_result_child3_father == 6 &&
+ "Reference Child_3 Father Foo failed");
+ assert(ref_result_child3_as_mother == 5 &&
+ "Reference Child_3 Mother Parent Cast Foo failed");
+ assert(ref_result_child3_as_father == 6 &&
+ "Reference Child_3 Father Parent Cast Foo failed");
+
+ return 0;
+}
+
+int test_multiple_inheritance_implicit() {
+ Mother mother;
+ Father father;
+ Child_1 child_1;
+ Child_2 child_2;
+ Child_3 child_3;
+
+ // map results back to host
+ int result_mother, result_father;
+ int result_child1_father, result_child1_mother, result_child1_as_mother,
+ result_child1_as_father;
+ int result_child2_mother, result_child2_father, result_child2_as_mother,
+ result_child2_as_father;
+ int result_child3_mother, result_child3_father, result_child3_as_mother,
+ result_child3_as_father;
+
+ // Add reference-based results
+ int ref_result_mother, ref_result_father;
+ int ref_result_child1_father, ref_result_child1_mother,
+ ref_result_child1_as_mother, ref_result_child1_as_father;
+ int ref_result_child2_mother, ref_result_child2_father,
+ ref_result_child2_as_mother, ref_result_child2_as_father;
+ int ref_result_child3_mother, ref_result_child3_father,
+ ref_result_child3_as_mother, ref_result_child3_as_father;
+
+#pragma omp target data map(father, mother, child_1, child_2, child_3)
+ {
+ // Base class pointers and references
+ Mother *ptr_mother = &mother;
+ Father *ptr_father = &father;
+ Mother &ref_mother = mother;
+ Father &ref_father = father;
+
+ // Child_1 pointers, references and casts
+ Child_1 *ptr_child_1 = &child_1;
+ Mother *ptr_child_1_cast_mother = &child_1;
+ Father *ptr_child_1_cast_father = &child_1;
+ Child_1 &ref_child_1 = child_1;
+ Mother &ref_child_1_cast_mother = child_1;
+ Father &ref_child_1_cast_father = child_1;
+
+ // Child_2 pointers, references and casts
+ Child_2 *ptr_child_2 = &child_2;
+ Mother *ptr_child_2_cast_mother = &child_2;
+ Father *ptr_child_2_cast_father = &child_2;
+ Child_2 &ref_child_2 = child_2;
+ Mother &ref_child_2_cast_mother = child_2;
+ Father &ref_child_2_cast_father = child_2;
+
+ // Child_3 pointers and casts
+ Child_3 *ptr_child_3 = &child_3;
+ Mother *ptr_child_3_cast_mother = &child_3;
+ Father *ptr_child_3_cast_father = &child_3;
+ Child_3 &ref_child_3 = child_3;
+ Mother &ref_child_3_cast_mother = child_3;
+ Father &ref_child_3_cast_father = child_3;
+
+ // Implicit mapping test - no explicit map clauses for pointers/references
+#pragma omp target map( \
+ from : result_mother, result_father, result_child1_father, \
+ result_child1_mother, result_child1_as_mother, \
+ result_child1_as_father, result_child2_mother, \
+ result_child2_father, result_child2_as_mother, \
+ result_child2_as_father, result_child3_mother, \
+ result_child3_father, result_child3_as_mother, \
+ result_child3_as_father, ref_result_mother, ref_result_father, \
+ ref_result_child1_father, ref_result_child1_mother, \
+ ref_result_child1_as_mother, ref_result_child1_as_father, \
+ ref_result_child2_mother, ref_result_child2_father, \
+ ref_result_child2_as_mother, ref_result_child2_as_father, \
+ ref_result_child3_mother, ref_result_child3_father, \
+ ref_result_child3_as_mother, ref_result_child3_as_father)
+ {
+ // These calls will fail if Clang does not
+ // translate/attach the vtable pointer in each object
+
+ // Pointer-based calls
+ // Mother
+ result_mother = ptr_mother->MotherFoo(1);
+ // Father
+ result_father = ptr_father->FatherFoo(1);
+ // Child_1
+ result_child1_father = ptr_child_1->FatherFoo(1);
+ result_child1_mother = ptr_child_1->MotherFoo(1);
+ result_child1_as_mother = ptr_child_1_cast_mother->MotherFoo(1);
+ result_child1_as_father = ptr_child_1_cast_father->FatherFoo(1);
+ // Child_2
+ result_child2_mother = ptr_child_2->MotherFoo(1);
+ result_child2_father = ptr_child_2->FatherFoo(1);
+ result_child2_as_mother = ptr_child_2_cast_mother->MotherFoo(1);
+ result_child2_as_father = ptr_child_2_cast_father->FatherFoo(1);
+ // Child_3
+ result_child3_mother = ptr_child_3->MotherFoo(1);
+ result_child3_father = ptr_child_3->FatherFoo(1);
+ result_child3_as_mother = ptr_child_3_cast_mother->MotherFoo(1);
+ result_child3_as_father = ptr_child_3_cast_father->FatherFoo(1);
+
+ // Reference-based calls
+ // Mother
+ ref_result_mother = ref_mother.MotherFoo(1);
+ // Father
+ ref_result_father = ref_father.FatherFoo(1);
+ // Child_1
+ ref_result_child1_father = ref_child_1.FatherFoo(1);
+ ref_result_child1_mother = ref_child_1.MotherFoo(1);
+ ref_result_child1_as_mother = ref_child_1_cast_mother.MotherFoo(1);
+ ref_result_child1_as_father = ref_child_1_cast_father.FatherFoo(1);
+ // Child_2
+ ref_result_child2_mother = ref_child_2.MotherFoo(1);
+ ref_result_child2_father = ref_child_2.FatherFoo(1);
+ ref_result_child2_as_mother = ref_child_2_cast_mother.MotherFoo(1);
+ ref_result_child2_as_father = ref_child_2_cast_father.FatherFoo(1);
+ // Child_3
+ ref_result_child3_mother = ref_child_3.MotherFoo(1);
+ ref_result_child3_father = ref_child_3.FatherFoo(1);
+ ref_result_child3_as_mother = ref_child_3_cast_mother.MotherFoo(1);
+ ref_result_child3_as_father = ref_child_3_cast_father.FatherFoo(1);
+ }
+ }
+
+ // Check pointer-based results
+ assert(result_mother == 1 && "Implicit Mother Foo failed");
+ assert(result_father == 2 && "Implicit Father Foo failed");
+ assert(result_child1_father == 3 && "Implicit Child_1 Father Foo failed");
+ assert(result_child1_mother == 1 && "Implicit Child_1 Mother Foo failed");
+ assert(result_child1_as_mother == 1 &&
+ "Implicit Child_1 Mother Parent Cast Foo failed");
+ assert(result_child1_as_father == 3 &&
+ "Implicit Child_1 Father Parent Cast Foo failed");
+ assert(result_child2_mother == 4 && "Implicit Child_2 Mother Foo failed");
+ assert(result_child2_father == 2 && "Implicit Child_2 Father Foo failed");
+ assert(result_child2_as_mother == 4 &&
+ "Implicit Child_2 Mother Parent Cast Foo failed");
+ assert(result_child2_as_father == 2 &&
+ "Implicit Child_2 Father Parent Cast Foo failed");
+ assert(result_child3_mother == 5 && "Implicit Child_3 Mother Foo failed");
+ assert(result_child3_father == 6 && "Implicit Child_3 Father Foo failed");
+ assert(result_child3_as_mother == 5 &&
+ "Implicit Child_3 Mother Parent Cast Foo failed");
+ assert(result_child3_as_father == 6 &&
+ "Implicit Child_3 Father Parent Cast Foo failed");
+
+ // Check reference-based results
+ assert(ref_result_mother == 1 && "Implicit Reference Mother Foo failed");
+ assert(ref_result_father == 2 && "Implicit Reference Father Foo failed");
+ assert(ref_result_child1_father == 3 &&
+ "Implicit Reference Child_1 Father Foo failed");
+ assert(ref_result_child1_mother == 1 &&
+ "Implicit Reference Child_1 Mother Foo failed");
+ assert(ref_result_child1_as_mother == 1 &&
+ "Implicit Reference Child_1 Mother Parent Cast Foo failed");
+ assert(ref_result_child1_as_father == 3 &&
+ "Implicit Reference Child_1 Father Parent Cast Foo failed");
+ assert(ref_result_child2_mother == 4 &&
+ "Implicit Reference Child_2 Mother Foo failed");
+ assert(ref_result_child2_father == 2 &&
+ "Implicit Reference Child_2 Father Foo failed");
+ assert(ref_result_child2_as_mother == 4 &&
+ "Implicit Reference Child_2 Mother Parent Cast Foo failed");
+ assert(ref_result_child2_as_father == 2 &&
+ "Implicit Reference Child_2 Father Parent Cast Foo failed");
+ assert(ref_result_child3_mother == 5 &&
+ "Implicit Reference Child_3 Mother Foo failed");
+ assert(ref_result_child3_father == 6 &&
+ "Implicit Reference Child_3 Father Foo failed");
+ assert(ref_result_child3_as_mother == 5 &&
+ "Implicit Reference Child_3 Mother Parent Cast Foo failed");
+ assert(ref_result_child3_as_father == 6 &&
+ "Implicit Reference Child_3 Father Parent Cast Foo failed");
+
+ return 0;
+}
+
+int main() {
+ test_multiple_inheritance();
+ test_multiple_inheritance_implicit();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
new file mode 100644
index 0000000000000..8a716bcf679ef
--- /dev/null
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
@@ -0,0 +1,428 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp declare target
+
+class Parent1 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ Parent1Foo(int x) {
+ return x;
+ }
+};
+
+class Parent2 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ Parent2Foo(int x) {
+ return 2 * x;
+ }
+};
+
+class Parent3 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ Parent3Foo(int x) {
+ return 3 * x;
+ }
+};
+
+class Parent4 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ Parent4Foo(int x) {
+ return 4 * x;
+ }
+};
+
+class Parent5 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ Parent5Foo(int x) {
+ return 5 * x;
+ }
+};
+
+class Child : public Parent1,
+ public Parent2,
+ public Parent3,
+ public Parent4,
+ public Parent5 {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) int
+ Parent1Foo(int x) override {
+ return 6 * x;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) int
+ Parent2Foo(int x) override {
+ return 7 * x;
+ }
+ __attribute__((noinline)) __attribute__((optnone)) int
+ Parent3Foo(int x) override {
+ return 8 * x;
+ }
+
+ // parent 4 stays the same
+
+ __attribute__((noinline)) __attribute__((optnone)) int
+ Parent5Foo(int x) override {
+ return 10 * x;
+ }
+};
+
+#pragma omp end declare target
+
+int test_five_parent_inheritance() {
+ Parent1 parent1;
+ Parent2 parent2;
+ Parent3 parent3;
+ Parent4 parent4;
+ Parent5 parent5;
+ Child child;
+
+ // map results back to host
+ int result_parent1, result_parent2, result_parent3, result_parent4,
+ result_parent5;
+ int result_child_parent1, result_child_parent2, result_child_parent3,
+ result_child_parent4, result_child_parent5;
+ int result_child_as_parent1, result_child_as_parent2, result_child_as_parent3,
+ result_child_as_parent4, result_child_as_parent5;
+
+ // Add reference-based results
+ int ref_result_parent1, ref_result_parent2, ref_result_parent3,
+ ref_result_parent4, ref_result_parent5;
+ int ref_result_child_parent1, ref_result_child_parent2,
+ ref_result_child_parent3, ref_result_child_parent4,
+ ref_result_child_parent5;
+ int ref_result_child_as_parent1, ref_result_child_as_parent2,
+ ref_result_child_as_parent3, ref_result_child_as_parent4,
+ ref_result_child_as_parent5;
+
+#pragma omp target data map(parent1, parent2, parent3, parent4, parent5, child)
+ {
+ // Base class pointers
+ Parent1 *ptr_parent1 = &parent1;
+ Parent2 *ptr_parent2 = &parent2;
+ Parent3 *ptr_parent3 = &parent3;
+ Parent4 *ptr_parent4 = &parent4;
+ Parent5 *ptr_parent5 = &parent5;
+
+ // Base class references
+ Parent1 &ref_parent1 = parent1;
+ Parent2 &ref_parent2 = parent2;
+ Parent3 &ref_parent3 = parent3;
+ Parent4 &ref_parent4 = parent4;
+ Parent5 &ref_parent5 = parent5;
+
+ // Child pointers
+ Child *ptr_child = &child;
+ Parent1 *ptr_child_cast_parent1 = &child;
+ Parent2 *ptr_child_cast_parent2 = &child;
+ Parent3 *ptr_child_cast_parent3 = &child;
+ Parent4 *ptr_child_cast_parent4 = &child;
+ Parent5 *ptr_child_cast_parent5 = &child;
+
+ // Child references
+ Child &ref_child = child;
+ Parent1 &ref_child_cast_parent1 = child;
+ Parent2 &ref_child_cast_parent2 = child;
+ Parent3 &ref_child_cast_parent3 = child;
+ Parent4 &ref_child_cast_parent4 = child;
+ Parent5 &ref_child_cast_parent5 = child;
+
+#pragma omp target map( \
+ from : result_parent1, result_parent2, result_parent3, result_parent4, \
+ result_parent5, result_child_parent1, result_child_parent2, \
+ result_child_parent3, result_child_parent4, result_child_parent5, \
+ result_child_as_parent1, result_child_as_parent2, \
+ result_child_as_parent3, result_child_as_parent4, \
+ result_child_as_parent5, ref_result_parent1, ref_result_parent2, \
+ ref_result_parent3, ref_result_parent4, ref_result_parent5, \
+ ref_result_child_parent1, ref_result_child_parent2, \
+ ref_result_child_parent3, ref_result_child_parent4, \
+ ref_result_child_parent5, ref_result_child_as_parent1, \
+ ref_result_child_as_parent2, ref_result_child_as_parent3, \
+ ref_result_child_as_parent4, ref_result_child_as_parent5) \
+ map(ptr_parent1[0 : 0], ptr_parent2[0 : 0], ptr_parent3[0 : 0], \
+ ptr_parent4[0 : 0], ptr_parent5[0 : 0], ptr_child[0 : 0], \
+ ptr_child_cast_parent1[0 : 0], ptr_child_cast_parent2[0 : 0], \
+ ptr_child_cast_parent3[0 : 0], ptr_child_cast_parent4[0 : 0], \
+ ptr_child_cast_parent5[0 : 0], ref_parent1, ref_parent2, \
+ ref_parent3, ref_parent4, ref_parent5, ref_child, \
+ ref_child_cast_parent1, ref_child_cast_parent2, \
+ ref_child_cast_parent3, ref_child_cast_parent4, \
+ ref_child_cast_parent5)
+ {
+ // Base class calls using pointers
+ result_parent1 = ptr_parent1->Parent1Foo(1);
+ result_parent2 = ptr_parent2->Parent2Foo(1);
+ result_parent3 = ptr_parent3->Parent3Foo(1);
+ result_parent4 = ptr_parent4->Parent4Foo(1);
+ result_parent5 = ptr_parent5->Parent5Foo(1);
+
+ // Direct child calls using pointers
+ result_child_parent1 = ptr_child->Parent1Foo(1);
+ result_child_parent2 = ptr_child->Parent2Foo(1);
+ result_child_parent3 = ptr_child->Parent3Foo(1);
+ result_child_parent4 = ptr_child->Parent4Foo(1);
+ result_child_parent5 = ptr_child->Parent5Foo(1);
+
+ // Polymorphic calls through parent pointers
+ result_child_as_parent1 = ptr_child_cast_parent1->Parent1Foo(1);
+ result_child_as_parent2 = ptr_child_cast_parent2->Parent2Foo(1);
+ result_child_as_parent3 = ptr_child_cast_parent3->Parent3Foo(1);
+ result_child_as_parent4 = ptr_child_cast_parent4->Parent4Foo(1);
+ result_child_as_parent5 = ptr_child_cast_parent5->Parent5Foo(1);
+
+ // Base class calls using references
+ ref_result_parent1 = ref_parent1.Parent1Foo(1);
+ ref_result_parent2 = ref_parent2.Parent2Foo(1);
+ ref_result_parent3 = ref_parent3.Parent3Foo(1);
+ ref_result_parent4 = ref_parent4.Parent4Foo(1);
+ ref_result_parent5 = ref_parent5.Parent5Foo(1);
+
+ // Direct child calls using references
+ ref_result_child_parent1 = ref_child.Parent1Foo(1);
+ ref_result_child_parent2 = ref_child.Parent2Foo(1);
+ ref_result_child_parent3 = ref_child.Parent3Foo(1);
+ ref_result_child_parent4 = ref_child.Parent4Foo(1);
+ ref_result_child_parent5 = ref_child.Parent5Foo(1);
+
+ // Polymorphic calls through parent references
+ ref_result_child_as_parent1 = ref_child_cast_parent1.Parent1Foo(1);
+ ref_result_child_as_parent2 = ref_child_cast_parent2.Parent2Foo(1);
+ ref_result_child_as_parent3 = ref_child_cast_parent3.Parent3Foo(1);
+ ref_result_child_as_parent4 = ref_child_cast_parent4.Parent4Foo(1);
+ ref_result_child_as_parent5 = ref_child_cast_parent5.Parent5Foo(1);
+ }
+ }
+
+ // Verify pointer-based results
+ assert(result_parent1 == 1 && "Parent1 Foo failed");
+ assert(result_parent2 == 2 && "Parent2 Foo failed");
+ assert(result_parent3 == 3 && "Parent3 Foo failed");
+ assert(result_parent4 == 4 && "Parent4 Foo failed");
+ assert(result_parent5 == 5 && "Parent5 Foo failed");
+
+ assert(result_child_parent1 == 6 && "Child Parent1 Foo failed");
+ assert(result_child_parent2 == 7 && "Child Parent2 Foo failed");
+ assert(result_child_parent3 == 8 && "Child Parent3 Foo failed");
+ assert(result_child_parent4 == 4 && "Child Parent4 Foo failed");
+ assert(result_child_parent5 == 10 && "Child Parent5 Foo failed");
+
+ assert(result_child_as_parent1 == 6 && "Child Parent1 Cast Foo failed");
+ assert(result_child_as_parent2 == 7 && "Child Parent2 Cast Foo failed");
+ assert(result_child_as_parent3 == 8 && "Child Parent3 Cast Foo failed");
+ assert(result_child_as_parent4 == 4 && "Child Parent4 Cast Foo failed");
+ assert(result_child_as_parent5 == 10 && "Child Parent5 Cast Foo failed");
+
+ // Verify reference-based results
+ assert(ref_result_parent1 == 1 && "Reference Parent1 Foo failed");
+ assert(ref_result_parent2 == 2 && "Reference Parent2 Foo failed");
+ assert(ref_result_parent3 == 3 && "Reference Parent3 Foo failed");
+ assert(ref_result_parent4 == 4 && "Reference Parent4 Foo failed");
+ assert(ref_result_parent5 == 5 && "Reference Parent5 Foo failed");
+
+ assert(ref_result_child_parent1 == 6 && "Reference Child Parent1 Foo failed");
+ assert(ref_result_child_parent2 == 7 && "Reference Child Parent2 Foo failed");
+ assert(ref_result_child_parent3 == 8 && "Reference Child Parent3 Foo failed");
+ assert(ref_result_child_parent4 == 4 && "Reference Child Parent4 Foo failed");
+ assert(ref_result_child_parent5 == 10 &&
+ "Reference Child Parent5 Foo failed");
+
+ assert(ref_result_child_as_parent1 == 6 &&
+ "Reference Child Parent1 Cast Foo failed");
+ assert(ref_result_child_as_parent2 == 7 &&
+ "Reference Child Parent2 Cast Foo failed");
+ assert(ref_result_child_as_parent3 == 8 &&
+ "Reference Child Parent3 Cast Foo failed");
+ assert(ref_result_child_as_parent4 == 4 &&
+ "Reference Child Parent4 Cast Foo failed");
+ assert(ref_result_child_as_parent5 == 10 &&
+ "Reference Child Parent5 Cast Foo failed");
+
+ return 0;
+}
+
+int test_five_parent_inheritance_implicit() {
+ Parent1 parent1;
+ Parent2 parent2;
+ Parent3 parent3;
+ Parent4 parent4;
+ Parent5 parent5;
+ Child child;
+
+ // map results back to host
+ int result_parent1, result_parent2, result_parent3, result_parent4,
+ result_parent5;
+ int result_child_parent1, result_child_parent2, result_child_parent3,
+ result_child_parent4, result_child_parent5;
+ int result_child_as_parent1, result_child_as_parent2, result_child_as_parent3,
+ result_child_as_parent4, result_child_as_parent5;
+
+ // Add reference-based results
+ int ref_result_parent1, ref_result_parent2, ref_result_parent3,
+ ref_result_parent4, ref_result_parent5;
+ int ref_result_child_parent1, ref_result_child_parent2,
+ ref_result_child_parent3, ref_result_child_parent4,
+ ref_result_child_parent5;
+ int ref_result_child_as_parent1, ref_result_child_as_parent2,
+ ref_result_child_as_parent3, ref_result_child_as_parent4,
+ ref_result_child_as_parent5;
+
+#pragma omp target data map(parent1, parent2, parent3, parent4, parent5, child)
+ {
+ // Base class pointers
+ Parent1 *ptr_parent1 = &parent1;
+ Parent2 *ptr_parent2 = &parent2;
+ Parent3 *ptr_parent3 = &parent3;
+ Parent4 *ptr_parent4 = &parent4;
+ Parent5 *ptr_parent5 = &parent5;
+
+ // Base class references
+ Parent1 &ref_parent1 = parent1;
+ Parent2 &ref_parent2 = parent2;
+ Parent3 &ref_parent3 = parent3;
+ Parent4 &ref_parent4 = parent4;
+ Parent5 &ref_parent5 = parent5;
+
+ // Child pointers
+ Child *ptr_child = &child;
+ Parent1 *ptr_child_cast_parent1 = &child;
+ Parent2 *ptr_child_cast_parent2 = &child;
+ Parent3 *ptr_child_cast_parent3 = &child;
+ Parent4 *ptr_child_cast_parent4 = &child;
+ Parent5 *ptr_child_cast_parent5 = &child;
+
+ // Child references
+ Child &ref_child = child;
+ Parent1 &ref_child_cast_parent1 = child;
+ Parent2 &ref_child_cast_parent2 = child;
+ Parent3 &ref_child_cast_parent3 = child;
+ Parent4 &ref_child_cast_parent4 = child;
+ Parent5 &ref_child_cast_parent5 = child;
+
+#pragma omp target map( \
+ from : result_parent1, result_parent2, result_parent3, result_parent4, \
+ result_parent5, result_child_parent1, result_child_parent2, \
+ result_child_parent3, result_child_parent4, result_child_parent5, \
+ result_child_as_parent1, result_child_as_parent2, \
+ result_child_as_parent3, result_child_as_parent4, \
+ result_child_as_parent5, ref_result_parent1, ref_result_parent2, \
+ ref_result_parent3, ref_result_parent4, ref_result_parent5, \
+ ref_result_child_parent1, ref_result_child_parent2, \
+ ref_result_child_parent3, ref_result_child_parent4, \
+ ref_result_child_parent5, ref_result_child_as_parent1, \
+ ref_result_child_as_parent2, ref_result_child_as_parent3, \
+ ref_result_child_as_parent4, ref_result_child_as_parent5)
+ {
+ // Base class calls using pointers
+ result_parent1 = ptr_parent1->Parent1Foo(1);
+ result_parent2 = ptr_parent2->Parent2Foo(1);
+ result_parent3 = ptr_parent3->Parent3Foo(1);
+ result_parent4 = ptr_parent4->Parent4Foo(1);
+ result_parent5 = ptr_parent5->Parent5Foo(1);
+
+ // Direct child calls using pointers
+ result_child_parent1 = ptr_child->Parent1Foo(1);
+ result_child_parent2 = ptr_child->Parent2Foo(1);
+ result_child_parent3 = ptr_child->Parent3Foo(1);
+ result_child_parent4 = ptr_child->Parent4Foo(1);
+ result_child_parent5 = ptr_child->Parent5Foo(1);
+
+ // Polymorphic calls through parent pointers
+ result_child_as_parent1 = ptr_child_cast_parent1->Parent1Foo(1);
+ result_child_as_parent2 = ptr_child_cast_parent2->Parent2Foo(1);
+ result_child_as_parent3 = ptr_child_cast_parent3->Parent3Foo(1);
+ result_child_as_parent4 = ptr_child_cast_parent4->Parent4Foo(1);
+ result_child_as_parent5 = ptr_child_cast_parent5->Parent5Foo(1);
+
+ // Base class calls using references
+ ref_result_parent1 = ref_parent1.Parent1Foo(1);
+ ref_result_parent2 = ref_parent2.Parent2Foo(1);
+ ref_result_parent3 = ref_parent3.Parent3Foo(1);
+ ref_result_parent4 = ref_parent4.Parent4Foo(1);
+ ref_result_parent5 = ref_parent5.Parent5Foo(1);
+
+ // Direct child calls using references
+ ref_result_child_parent1 = ref_child.Parent1Foo(1);
+ ref_result_child_parent2 = ref_child.Parent2Foo(1);
+ ref_result_child_parent3 = ref_child.Parent3Foo(1);
+ ref_result_child_parent4 = ref_child.Parent4Foo(1);
+ ref_result_child_parent5 = ref_child.Parent5Foo(1);
+
+ // Polymorphic calls through parent references
+ ref_result_child_as_parent1 = ref_child_cast_parent1.Parent1Foo(1);
+ ref_result_child_as_parent2 = ref_child_cast_parent2.Parent2Foo(1);
+ ref_result_child_as_parent3 = ref_child_cast_parent3.Parent3Foo(1);
+ ref_result_child_as_parent4 = ref_child_cast_parent4.Parent4Foo(1);
+ ref_result_child_as_parent5 = ref_child_cast_parent5.Parent5Foo(1);
+ }
+ }
+ // Verify pointer-based results
+ assert(result_parent1 == 1 && "Implicit Parent1 Foo failed");
+ assert(result_parent2 == 2 && "Implicit Parent2 Foo failed");
+ assert(result_parent3 == 3 && "Implicit Parent3 Foo failed");
+ assert(result_parent4 == 4 && "Implicit Parent4 Foo failed");
+ assert(result_parent5 == 5 && "Implicit Parent5 Foo failed");
+
+ assert(result_child_parent1 == 6 && "Implicit Child Parent1 Foo failed");
+ assert(result_child_parent2 == 7 && "Implicit Child Parent2 Foo failed");
+ assert(result_child_parent3 == 8 && "Implicit Child Parent3 Foo failed");
+ assert(result_child_parent4 == 4 && "Implicit Child Parent4 Foo failed");
+ assert(result_child_parent5 == 10 && "Implicit Child Parent5 Foo failed");
+
+ assert(result_child_as_parent1 == 6 &&
+ "Implicit Child Parent1 Cast Foo failed");
+ assert(result_child_as_parent2 == 7 &&
+ "Implicit Child Parent2 Cast Foo failed");
+ assert(result_child_as_parent3 == 8 &&
+ "Implicit Child Parent3 Cast Foo failed");
+ assert(result_child_as_parent4 == 4 &&
+ "Implicit Child Parent4 Cast Foo failed");
+ assert(result_child_as_parent5 == 10 &&
+ "Implicit Child Parent5 Cast Foo failed");
+
+ // Verify reference-based results
+ assert(ref_result_parent1 == 1 && "Implicit Reference Parent1 Foo failed");
+ assert(ref_result_parent2 == 2 && "Implicit Reference Parent2 Foo failed");
+ assert(ref_result_parent3 == 3 && "Implicit Reference Parent3 Foo failed");
+ assert(ref_result_parent4 == 4 && "Implicit Reference Parent4 Foo failed");
+ assert(ref_result_parent5 == 5 && "Implicit Reference Parent5 Foo failed");
+
+ assert(ref_result_child_parent1 == 6 &&
+ "Implicit Reference Child Parent1 Foo failed");
+ assert(ref_result_child_parent2 == 7 &&
+ "Implicit Reference Child Parent2 Foo failed");
+ assert(ref_result_child_parent3 == 8 &&
+ "Implicit Reference Child Parent3 Foo failed");
+ assert(ref_result_child_parent4 == 4 &&
+ "Implicit Reference Child Parent4 Foo failed");
+ assert(ref_result_child_parent5 == 10 &&
+ "Implicit Reference Child Parent5 Foo failed");
+
+ assert(ref_result_child_as_parent1 == 6 &&
+ "Implicit Reference Child Parent1 Cast Foo failed");
+ assert(ref_result_child_as_parent2 == 7 &&
+ "Implicit Reference Child Parent2 Cast Foo failed");
+ assert(ref_result_child_as_parent3 == 8 &&
+ "Implicit Reference Child Parent3 Cast Foo failed");
+ assert(ref_result_child_as_parent4 == 4 &&
+ "Implicit Reference Child Parent4 Cast Foo failed");
+ assert(ref_result_child_as_parent5 == 10 &&
+ "Implicit Reference Child Parent5 Cast Foo failed");
+
+ return 0;
+}
+
+int main() {
+ test_five_parent_inheritance();
+ test_five_parent_inheritance_implicit();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_virtual_func_reference.cpp b/offload/test/api/omp_virtual_func_reference.cpp
new file mode 100644
index 0000000000000..47930d974f0a7
--- /dev/null
+++ b/offload/test/api/omp_virtual_func_reference.cpp
@@ -0,0 +1,80 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#define TEST_VAL 10
+
+#pragma omp declare target
+class Base {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo(int x) {
+ return x;
+ }
+};
+
+class Derived : public Base {
+public:
+ __attribute__((noinline)) __attribute__((optnone)) virtual int foo(int x) {
+ return -x;
+ }
+};
+#pragma omp end declare target
+
+int test_virtual_reference() {
+ Derived ddd;
+ Base cont;
+ Base &bbb = ddd;
+
+ int b_ret, d_ret, c_ret;
+
+#pragma omp target data map(to : ddd, cont)
+ {
+#pragma omp target map(bbb, ddd, cont) map(from : b_ret, d_ret, c_ret)
+ {
+ b_ret = bbb.foo(TEST_VAL);
+ d_ret = ddd.foo(TEST_VAL);
+ c_ret = cont.foo(TEST_VAL);
+ }
+ }
+
+ assert(c_ret == TEST_VAL && "Control Base call failed on gpu");
+ assert(b_ret == -TEST_VAL && "Control Base call failed on gpu");
+ assert(d_ret == -TEST_VAL && "Derived call failed on gpu");
+
+ return 0;
+}
+
+int test_virtual_reference_implicit() {
+ Derived ddd;
+ Base cont;
+ Base &bbb = ddd;
+
+ int b_ret, d_ret, c_ret;
+
+#pragma omp target data map(to : ddd, cont)
+ {
+#pragma omp target map(from : b_ret, d_ret, c_ret)
+ {
+ b_ret = bbb.foo(TEST_VAL);
+ d_ret = ddd.foo(TEST_VAL);
+ c_ret = cont.foo(TEST_VAL);
+ }
+ }
+
+ assert(c_ret == TEST_VAL && "Control Base call failed on gpu");
+ assert(b_ret == -TEST_VAL && "Control Base call failed on gpu");
+ assert(d_ret == -TEST_VAL && "Derived call failed on gpu");
+
+ return 0;
+}
+
+int main() {
+ test_virtual_reference();
+ test_virtual_reference_implicit();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
More information about the llvm-commits
mailing list