[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
Tue Jan 6 08:24:02 PST 2026
https://github.com/Jason-VanBeusekom updated https://github.com/llvm/llvm-project/pull/159857
>From 37847a2e289dcd687a47156a6779c781e3da77d3 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/4] [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/CodeGenModule.h | 3 +
clang/test/OpenMP/target_vtable_codegen.cpp | 280 ++++++++++++++++++++
2 files changed, 283 insertions(+)
create mode 100644 clang/test/OpenMP/target_vtable_codegen.cpp
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 38b052e5cd1dd..18a8855ae443a 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -783,6 +783,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
>From 8e96b16b5ea00f7a75d4e23a76018eaa70bba378 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/4] [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 | 29 ++
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, 1513 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 ba2cdd3ea19dc..0e35615643a3c 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6824,6 +6824,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 01661ad54ee2f..bef0e86c7b627 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6342,6 +6342,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);
@@ -6350,6 +6369,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 a6c80cd083bb8..6cea3b87e45dc 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -2271,6 +2271,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 d4ed7bdb472cb59714c18b241814b73495587efa 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/4] 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 0e35615643a3c..d9b030f811c4a 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6825,17 +6825,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 bef0e86c7b627..2a91e1c90ca5d 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6350,9 +6350,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 6cea3b87e45dc..c9bc086da12d1 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -2271,17 +2271,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 5d5a2a383f2b2..bcc35558ca101 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);
}
>From 1619e73d908edb54ca2a4121f4faf40ca5f06d84 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 6 Jan 2026 10:08:22 -0600
Subject: [PATCH 4/4] Revert "[OpenMP][clang] Register Vtables on device for
indirect calls"
This reverts commit 77fd376b5b87eab76bda14e7e457ea80e8e09f20.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 11 +-
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 | 4 +-
7 files changed, 8 insertions(+), 305 deletions(-)
delete mode 100644 clang/test/OpenMP/target_vtable_codegen.cpp
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 2a91e1c90ca5d..92c0e7785c54c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1776,16 +1776,13 @@ 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);
}
+<<<<<<< HEAD
void CGOpenMPRuntime::registerVTableOffloadEntry(llvm::GlobalVariable *VTable,
const VarDecl *VD) {
// TODO: add logic to avoid duplicate vtable registrations per
@@ -1896,6 +1893,8 @@ void CGOpenMPRuntime::registerVTable(const OMPExecutableDirective &D) {
}
}
+=======
+>>>>>>> be29867696e3 (Revert "[OpenMP][clang] Register Vtables on device for indirect calls")
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
QualType VarType,
StringRef Name) {
@@ -6398,7 +6397,6 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr);
}
}
- registerVTable(D);
}
/// Checks if the expression is constant or does not have non-trivial function
@@ -11053,6 +11051,7 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
if (!S)
return;
+<<<<<<< HEAD
// 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).
@@ -11064,6 +11063,8 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
registerVTable(*E);
}
+=======
+>>>>>>> be29867696e3 (Revert "[OpenMP][clang] Register Vtables on device for indirect calls")
// Codegen OMP target directives that offload compute to the device.
bool RequiresDeviceCodegen =
isa<OMPExecutableDirective>(S) &&
diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 46f00192c5735..f0d7e7003f2d9 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -38,12 +38,6 @@ 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 37458eee02e34..5c45e355fb145 100644
--- a/clang/lib/CodeGen/CGVTables.h
+++ b/clang/lib/CodeGen/CGVTables.h
@@ -122,10 +122,6 @@ 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 18a8855ae443a..38b052e5cd1dd 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -783,9 +783,6 @@ 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
deleted file mode 100644
index 276cef4eb8801..0000000000000
--- a/clang/test/OpenMP/target_vtable_codegen.cpp
+++ /dev/null
@@ -1,280 +0,0 @@
-///==========================================================================///
-// 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 c8db40d3cf51b..05d8a7dd168a3 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -396,8 +396,6 @@ 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
@@ -2763,8 +2761,7 @@ class OpenMPIRBuilder {
enum EmitMetadataErrorKind {
EMIT_MD_TARGET_REGION_ERROR,
EMIT_MD_DECLARE_TARGET_ERROR,
- EMIT_MD_GLOBAL_VAR_LINK_ERROR,
- EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR
+ EMIT_MD_GLOBAL_VAR_LINK_ERROR
};
/// Callback function type
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index b6a3d9e66fb9c..37e03f1dec655 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10658,9 +10658,7 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
// Indirect globals need to use a special name that doesn't match the name
// of the associated host global.
- if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect ||
- Flags ==
- OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable)
+ if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect)
createOffloadEntry(CE->getAddress(), CE->getAddress(), CE->getVarSize(),
Flags, CE->getLinkage(), CE->getVarName());
else
More information about the Openmp-commits
mailing list