[Openmp-commits] [clang] [llvm] [openmp] [OpenMP][clang] Indirect and Virtual function call mapping from host to device (PR #159857)
via Openmp-commits
openmp-commits at lists.llvm.org
Wed Oct 1 11:18:15 PDT 2025
https://github.com/Jason-VanBeusekom updated https://github.com/llvm/llvm-project/pull/159857
>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/3] [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/3] [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;
+}
>From 11b1f086b43736f07dd23313c474fab8a8e7e3e2 Mon Sep 17 00:00:00 2001
From: jason-van-beusekom <jason.van-beusekom at hpe.com>
Date: Wed, 1 Oct 2025 13:18:01 -0500
Subject: [PATCH 3/3] Updates based on feedback
---
clang/lib/CodeGen/CGExpr.cpp | 7 +++----
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 3 +--
clang/lib/CodeGen/ItaniumCXXABI.cpp | 11 +++++------
.../target_vtable_omp_indirect_call_lookup.cpp | 8 ++++----
offload/test/api/omp_indirect_call.c | 12 ++++++------
openmp/device/src/Misc.cpp | 2 +-
6 files changed, 20 insertions(+), 23 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index cc4c21a719f4c..15585ee7a829e 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6584,17 +6584,16 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
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
+ // Insert function pointer lookup if this is a target call
//
- // This is used for the indriect function Case, virtual function case is
+ // This is used for the indirect 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");
+ "__kmpc_omp_indirect_call_lookup");
llvm::Value *Func = Callee.getFunctionPointer();
llvm::Type *BackupTy = Func->getType();
Func = Builder.CreatePointerBitCastOrAddrSpaceCast(Func, PtrTy);
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index ac1d467affc00..01334ebd40e66 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6344,9 +6344,8 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
: CGF(CGF), TargetCalls(TargetCalls) {}
bool VisitCallExpr(CallExpr *CE) {
- if (!CE->getDirectCallee()) {
+ if (!CE->getDirectCallee())
TargetCalls.insert(CE);
- }
return true;
}
diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp
index 1dbfe23cef127..8937a3940fad1 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -2261,17 +2261,16 @@ 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.
- */
+
+ // For the translation of virtual functions, we need to map the (potential) host
+ // vtable to the device vtable. This is done by calling the runtime function
+ // __kmpc_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");
+ "__kmpc_omp_indirect_call_lookup");
auto *BackupTy = VTable->getType();
// Need to convert to generic address space
VTable = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(VTable, NewPtrTy);
diff --git a/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp b/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
index 52bbb382fb853..d9addd6291fcd 100644
--- a/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
+++ b/clang/test/OpenMP/target_vtable_omp_indirect_call_lookup.cpp
@@ -33,10 +33,10 @@ int main() {
#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]*}})
+ // CK1-DAG: call ptr @__kmpc_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__kmpc_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__kmpc_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
+ // CK1-DAG: call ptr @__kmpc_omp_indirect_call_lookup(ptr %vtable{{[0-9]*}})
int result1 = pointer1->foo();
int result2 = pointer1->bar();
int result3 = pointer2->foo();
diff --git a/offload/test/api/omp_indirect_call.c b/offload/test/api/omp_indirect_call.c
index ac0febf7854da..0484c8df0a33d 100644
--- a/offload/test/api/omp_indirect_call.c
+++ b/offload/test/api/omp_indirect_call.c
@@ -5,14 +5,14 @@
#pragma omp begin declare variant match(device = {kind(gpu)})
// Provided by the runtime.
-void *__llvm_omp_indirect_call_lookup(void *host_ptr);
-#pragma omp declare target to(__llvm_omp_indirect_call_lookup) \
+void *__kmpc_omp_indirect_call_lookup(void *host_ptr);
+#pragma omp declare target to(__kmpc_omp_indirect_call_lookup) \
device_type(nohost)
#pragma omp end declare variant
#pragma omp begin declare variant match(device = {kind(cpu)})
// We assume unified addressing on the CPU target.
-void *__llvm_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; }
+void *__kmpc_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; }
#pragma omp end declare variant
#pragma omp begin declare target indirect
@@ -32,11 +32,11 @@ int main() {
void *baz_res;
#pragma omp target map(to : foo_ptr, bar_ptr, baz_ptr) map(tofrom : count)
{
- foo_res = __llvm_omp_indirect_call_lookup(foo_ptr);
+ foo_res = __kmpc_omp_indirect_call_lookup(foo_ptr);
((void (*)(int *))foo_res)(&count);
- bar_res = __llvm_omp_indirect_call_lookup(bar_ptr);
+ bar_res = __kmpc_omp_indirect_call_lookup(bar_ptr);
((void (*)(int *))bar_res)(&count);
- baz_res = __llvm_omp_indirect_call_lookup(baz_ptr);
+ baz_res = __kmpc_omp_indirect_call_lookup(baz_ptr);
((void (*)(int *))baz_res)(&count);
}
diff --git a/openmp/device/src/Misc.cpp b/openmp/device/src/Misc.cpp
index a89f8b2a74531..a2383856a498e 100644
--- a/openmp/device/src/Misc.cpp
+++ b/openmp/device/src/Misc.cpp
@@ -89,7 +89,7 @@ double omp_get_wtime(void) {
return static_cast<double>(__builtin_readsteadycounter()) * omp_get_wtick();
}
-void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
+void *__kmpc_omp_indirect_call_lookup(void *HstPtr) {
return ompx::impl::indirectCallLookup(HstPtr);
}
More information about the Openmp-commits
mailing list