[clang] [llvm] [openmp] [OpenMP][clang] Indirect and Virtual function call mapping from host to device (PR #159857)
Jason Van Beusekom via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 2 07:34:39 PST 2026
https://github.com/Jason-Van-Beusekom updated https://github.com/llvm/llvm-project/pull/159857
>From f783b0889dc94fda28ac9086899332015c441db4 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 01/10] [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 b0960b472288683a599c4bae7e0034dce3318b7d 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 02/10] 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 746e8d88c8a38b3cb43c88d2d5fae7e73cb2d672 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 6 Jan 2026 11:46:30 -0600
Subject: [PATCH 03/10] fixes from rebase
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 1 +
clang/lib/CodeGen/CodeGenModule.h | 3 +++
2 files changed, 4 insertions(+)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 2a91e1c90ca5d..f732053a435a8 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/DiagnosticFrontend.h"
#include "clang/Basic/OpenMPKinds.h"
#include "clang/Basic/SourceManager.h"
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 38b052e5cd1dd..8fdf21edb2d24 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(); }
>From 553a35b8e81efa277ce1a5dfd5db138d94541f49 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 6 Jan 2026 15:57:34 -0600
Subject: [PATCH 04/10] format and name clean up
---
clang/lib/CodeGen/CGExpr.cpp | 2 +-
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 +-
clang/lib/CodeGen/ItaniumCXXABI.cpp | 11 +-
...target_vtable_omp_indirect_call_lookup.cpp | 8 +-
offload/test/api/omp_indirect_call.c | 12 +-
offload/test/api/omp_indirect_func_struct.c | 180 ++++++++++++------
offload/test/api/omp_virtual_func.cpp | 9 +-
openmp/device/src/Misc.cpp | 2 +-
8 files changed, 150 insertions(+), 78 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index d9b030f811c4a..37afb1aae6f9a 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6834,7 +6834,7 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
llvm::Type *RtlFnArgs[] = {PtrTy};
llvm::FunctionCallee DeviceRtlFn = CGM.CreateRuntimeFunction(
llvm::FunctionType::get(PtrTy, RtlFnArgs, false),
- "__kmpc_omp_indirect_call_lookup");
+ "__llvm_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 f732053a435a8..f398259cd94db 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -22,9 +22,9 @@
#include "clang/AST/Attr.h"
#include "clang/AST/Decl.h"
#include "clang/AST/OpenMPClause.h"
+#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/AST/StmtVisitor.h"
-#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/Basic/DiagnosticFrontend.h"
#include "clang/Basic/OpenMPKinds.h"
#include "clang/Basic/SourceManager.h"
@@ -6376,7 +6376,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
if (LangOpts.OpenMPIsTargetDevice) {
// Search AST for target "CallExpr"s of "OMPTargetAutoLookup".
OMPTargetCallCollector Visitor(CGF, CGF.CGM.OMPTargetCalls);
- Visitor.TraverseStmt(const_cast<Stmt*>(CS.getCapturedStmt()));
+ Visitor.TraverseStmt(const_cast<Stmt *>(CS.getCapturedStmt()));
}
CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp
index c9bc086da12d1..157e5ab2503bd 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -2271,16 +2271,17 @@ 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 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.
+
+ // 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
+ // __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),
- "__kmpc_omp_indirect_call_lookup");
+ "__llvm_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 d9addd6291fcd..52bbb382fb853 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 @__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]*}})
+ // 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();
diff --git a/offload/test/api/omp_indirect_call.c b/offload/test/api/omp_indirect_call.c
index 0484c8df0a33d..ac0febf7854da 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 *__kmpc_omp_indirect_call_lookup(void *host_ptr);
-#pragma omp declare target to(__kmpc_omp_indirect_call_lookup) \
+void *__llvm_omp_indirect_call_lookup(void *host_ptr);
+#pragma omp declare target to(__llvm_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 *__kmpc_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; }
+void *__llvm_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 = __kmpc_omp_indirect_call_lookup(foo_ptr);
+ foo_res = __llvm_omp_indirect_call_lookup(foo_ptr);
((void (*)(int *))foo_res)(&count);
- bar_res = __kmpc_omp_indirect_call_lookup(bar_ptr);
+ bar_res = __llvm_omp_indirect_call_lookup(bar_ptr);
((void (*)(int *))bar_res)(&count);
- baz_res = __kmpc_omp_indirect_call_lookup(baz_ptr);
+ baz_res = __llvm_omp_indirect_call_lookup(baz_ptr);
((void (*)(int *))baz_res)(&count);
}
diff --git a/offload/test/api/omp_indirect_func_struct.c b/offload/test/api/omp_indirect_func_struct.c
index cc2eeb86a2e5c..a21d7d3494ded 100644
--- a/offload/test/api/omp_indirect_func_struct.c
+++ b/offload/test/api/omp_indirect_func_struct.c
@@ -1,73 +1,101 @@
// RUN: %libomptarget-compile-run-and-check-generic
+#include <assert.h>
#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; }
+__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);
+ int (*indirect0)(int);
};
typedef struct {
int buffer;
int (*indirect1_ptr)();
- int (*indirect0_ptr)(int);
+ int (*indirect0_ptr)(int);
} indirect_stru_mapped;
-#pragma omp declare mapper (indirect_stru_mapped s) map(s,s.indirect0_ptr,s.indirect1_ptr)
+#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};
+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");
+ 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)
+#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");
+ 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)
+ 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");
+ 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() {
@@ -77,42 +105,58 @@ void test_local_struct_explicit_mapping() {
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");
+ 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)
+#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");
+ 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");
+ 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)
+#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");
+ 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() {
@@ -122,17 +166,23 @@ void test_global_struct_user_mapper() {
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");
+ 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)
+#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");
+ 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;
@@ -143,17 +193,23 @@ void test_global_struct_user_mapper() {
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");
+ 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)
+#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");
+ 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() {
@@ -167,17 +223,23 @@ void test_local_struct_user_mapper() {
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");
+ 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)
+#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");
+ 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;
@@ -188,17 +250,23 @@ void test_local_struct_user_mapper() {
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");
+ 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)
+#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");
+ 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() {
@@ -206,7 +274,7 @@ int main() {
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
index 1cfcb6f4d3a54..ba2e9b53b3686 100644
--- a/offload/test/api/omp_virtual_func.cpp
+++ b/offload/test/api/omp_virtual_func.cpp
@@ -15,7 +15,8 @@ class Base {
__attribute__((noinline)) __attribute__((optnone)) virtual int bar() {
return 2;
}
- __attribute__((noinline)) __attribute__((optnone)) virtual int foo_with_arg(int x) {
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ foo_with_arg(int x) {
return x;
}
};
@@ -28,7 +29,8 @@ class Derived : public Base {
__attribute__((noinline)) __attribute__((optnone)) virtual int bar() {
return 20;
}
- __attribute__((noinline)) __attribute__((optnone)) virtual int foo_with_arg(int x) {
+ __attribute__((noinline)) __attribute__((optnone)) virtual int
+ foo_with_arg(int x) {
return -x;
}
};
@@ -143,7 +145,8 @@ int test_virtual_reference_implicit() {
}
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(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;
diff --git a/openmp/device/src/Misc.cpp b/openmp/device/src/Misc.cpp
index bcc35558ca101..5d5a2a383f2b2 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 *__kmpc_omp_indirect_call_lookup(void *HstPtr) {
+void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
return ompx::impl::indirectCallLookup(HstPtr);
}
>From f6a09190efc1f964848523a5f0311aa4e041cf4a Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 6 Jan 2026 16:49:48 -0600
Subject: [PATCH 05/10] compile warning fix
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 8 +++-----
1 file changed, 3 insertions(+), 5 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index f398259cd94db..ecfde8a1b44f8 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6346,9 +6346,8 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
class OMPTargetCallCollector
: public RecursiveASTVisitor<OMPTargetCallCollector> {
public:
- OMPTargetCallCollector(CodeGenFunction &CGF,
- llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls)
- : CGF(CGF), TargetCalls(TargetCalls) {}
+ OMPTargetCallCollector(llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls)
+ : TargetCalls(TargetCalls) {}
bool VisitCallExpr(CallExpr *CE) {
if (!CE->getDirectCallee())
@@ -6357,7 +6356,6 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
}
private:
- CodeGenFunction &CGF;
llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls;
};
@@ -6375,7 +6373,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
const auto &LangOpts = CGF.getLangOpts();
if (LangOpts.OpenMPIsTargetDevice) {
// Search AST for target "CallExpr"s of "OMPTargetAutoLookup".
- OMPTargetCallCollector Visitor(CGF, CGF.CGM.OMPTargetCalls);
+ OMPTargetCallCollector Visitor(CGF.CGM.OMPTargetCalls);
Visitor.TraverseStmt(const_cast<Stmt *>(CS.getCapturedStmt()));
}
>From bb55fbb7474f87144719ed0ff765aacd68df51ac Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Wed, 7 Jan 2026 10:35:56 -0600
Subject: [PATCH 06/10] remove opnone and noinline attributes
---
offload/test/api/omp_indirect_func_basic.c | 8 +---
offload/test/api/omp_indirect_func_struct.c | 17 ++-----
offload/test/api/omp_virtual_func.cpp | 26 +++--------
...p_virtual_func_multiple_inheritance_01.cpp | 30 +++----------
...p_virtual_func_multiple_inheritance_02.cpp | 45 ++++---------------
.../test/api/omp_virtual_func_reference.cpp | 8 +---
6 files changed, 29 insertions(+), 105 deletions(-)
diff --git a/offload/test/api/omp_indirect_func_basic.c b/offload/test/api/omp_indirect_func_basic.c
index ff517247d4932..aaf0ccf371015 100644
--- a/offload/test/api/omp_indirect_func_basic.c
+++ b/offload/test/api/omp_indirect_func_basic.c
@@ -7,12 +7,8 @@
#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;
-}
+int direct(int x) { return 2 * x; }
+int indirect_base(int x) { return -1 * x; }
#pragma omp end declare target
int (*indirect)(int) = indirect_base;
diff --git a/offload/test/api/omp_indirect_func_struct.c b/offload/test/api/omp_indirect_func_struct.c
index a21d7d3494ded..f9e1489e4b4ea 100644
--- a/offload/test/api/omp_indirect_func_struct.c
+++ b/offload/test/api/omp_indirect_func_struct.c
@@ -7,19 +7,10 @@
#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;
-}
+int direct_arg(int x) { return 2 * x; }
+int indirect_base_arg(int x) { return -1 * x; }
+int direct() { return TEST_VAL; }
+int indirect_base() { return -1 * TEST_VAL; }
#pragma omp end declare target
struct indirect_stru {
diff --git a/offload/test/api/omp_virtual_func.cpp b/offload/test/api/omp_virtual_func.cpp
index ba2e9b53b3686..70d35e07cfbb3 100644
--- a/offload/test/api/omp_virtual_func.cpp
+++ b/offload/test/api/omp_virtual_func.cpp
@@ -9,30 +9,16 @@
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;
- }
+ virtual int foo() { return 1; }
+ virtual int bar() { return 2; }
+ 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;
- }
+ virtual int foo() { return 10; }
+ virtual int bar() { return 20; }
+ virtual int foo_with_arg(int x) { return -x; }
};
#pragma omp end declare target
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
index 20ab90cd35a3b..307ee78d01202 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
@@ -8,46 +8,28 @@
class Mother {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- MotherFoo(int x) {
- return x;
- }
+ virtual int MotherFoo(int x) { return x; }
};
class Father {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- FatherFoo(int x) {
- return x * 2;
- }
+ 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;
- }
+ 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;
- }
+ 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;
- }
+ virtual int MotherFoo(int x) { return x * 5; }
+ virtual int FatherFoo(int x) { return x * 6; }
};
#pragma omp end declare target
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
index 8a716bcf679ef..1aba0304d17e1 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
@@ -8,42 +8,27 @@
class Parent1 {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- Parent1Foo(int x) {
- return x;
- }
+ virtual int Parent1Foo(int x) { return x; }
};
class Parent2 {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- Parent2Foo(int x) {
- return 2 * x;
- }
+ virtual int Parent2Foo(int x) { return 2 * x; }
};
class Parent3 {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- Parent3Foo(int x) {
- return 3 * x;
- }
+ virtual int Parent3Foo(int x) { return 3 * x; }
};
class Parent4 {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- Parent4Foo(int x) {
- return 4 * x;
- }
+ virtual int Parent4Foo(int x) { return 4 * x; }
};
class Parent5 {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int
- Parent5Foo(int x) {
- return 5 * x;
- }
+ virtual int Parent5Foo(int x) { return 5 * x; }
};
class Child : public Parent1,
@@ -52,25 +37,13 @@ class Child : public Parent1,
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;
- }
+ int Parent1Foo(int x) override { return 6 * x; }
+ int Parent2Foo(int x) override { return 7 * x; }
+ 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;
- }
+ int Parent5Foo(int x) override { return 10 * x; }
};
#pragma omp end declare target
diff --git a/offload/test/api/omp_virtual_func_reference.cpp b/offload/test/api/omp_virtual_func_reference.cpp
index 47930d974f0a7..7fb25dbd5c626 100644
--- a/offload/test/api/omp_virtual_func_reference.cpp
+++ b/offload/test/api/omp_virtual_func_reference.cpp
@@ -9,16 +9,12 @@
#pragma omp declare target
class Base {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int foo(int x) {
- return x;
- }
+ virtual int foo(int x) { return x; }
};
class Derived : public Base {
public:
- __attribute__((noinline)) __attribute__((optnone)) virtual int foo(int x) {
- return -x;
- }
+ virtual int foo(int x) { return -x; }
};
#pragma omp end declare target
>From 7abe4866545d063d6478351c837d9e76f4684c1f Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 13 Jan 2026 16:50:37 -0600
Subject: [PATCH 07/10] Update indirect function lookup logic to be in Sema
---
clang/include/clang/AST/ASTContext.h | 4 ++++
clang/lib/CodeGen/CGExpr.cpp | 3 ++-
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 26 --------------------------
clang/lib/CodeGen/CodeGenModule.h | 3 ---
clang/lib/Sema/SemaOpenMP.cpp | 8 +++++++-
5 files changed, 13 insertions(+), 31 deletions(-)
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index 68205dd1c1fd9..dc27e2bd5d1d4 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -1370,6 +1370,10 @@ class ASTContext : public RefCountedBase<ASTContext> {
/// are stored here.
llvm::DenseMap<const CXXMethodDecl *, CXXCastPath> LambdaCastPaths;
+ /// Keep track of indirect call expressions within OpenMP target regions.
+ /// Used to instert calls to __llvm_omp_indirect_call_lookup during codegen.
+ llvm::DenseSet<const CallExpr *> OMPTargetCalls;
+
ASTContext(LangOptions &LOpts, SourceManager &SM, IdentifierTable &idents,
SelectorTable &sels, Builtin::Context &builtins,
TranslationUnitKind TUKind);
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 37afb1aae6f9a..8218be85484c6 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6829,7 +6829,8 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
//
// This is used for the indirect function case, virtual function case is
// handled in ItaniumCXXABI.cpp
- if (getLangOpts().OpenMPIsTargetDevice && CGM.OMPTargetCalls.contains(E)) {
+ if (getLangOpts().OpenMPIsTargetDevice &&
+ getContext().OMPTargetCalls.contains(E)) {
auto *PtrTy = CGM.VoidPtrTy;
llvm::Type *RtlFnArgs[] = {PtrTy};
llvm::FunctionCallee DeviceRtlFn = CGM.CreateRuntimeFunction(
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index ecfde8a1b44f8..cf99f7690312a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -22,7 +22,6 @@
#include "clang/AST/Attr.h"
#include "clang/AST/Decl.h"
#include "clang/AST/OpenMPClause.h"
-#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/AST/StmtVisitor.h"
#include "clang/Basic/DiagnosticFrontend.h"
@@ -6343,22 +6342,6 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
- class OMPTargetCallCollector
- : public RecursiveASTVisitor<OMPTargetCallCollector> {
- public:
- OMPTargetCallCollector(llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls)
- : TargetCalls(TargetCalls) {}
-
- bool VisitCallExpr(CallExpr *CE) {
- if (!CE->getDirectCallee())
- TargetCalls.insert(CE);
- return true;
- }
-
- private:
- llvm::SmallPtrSetImpl<const CallExpr *> &TargetCalls;
- };
-
llvm::TargetRegionEntryInfo EntryInfo =
getEntryInfoFromPresumedLoc(CGM, OMPBuilder, D.getBeginLoc(), ParentName);
@@ -6367,15 +6350,6 @@ 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.CGM.OMPTargetCalls);
- Visitor.TraverseStmt(const_cast<Stmt *>(CS.getCapturedStmt()));
- }
CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 8fdf21edb2d24..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/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 2a1337be13b99..b8711da4c26ea 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7360,8 +7360,14 @@ ExprResult SemaOpenMP::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
return Call;
FunctionDecl *CalleeFnDecl = CE->getDirectCallee();
- if (!CalleeFnDecl)
+
+ // Mark indirect calls inside target regions, to allow for insertion of
+ // __llvm_omp_indirect_call_lookup calls during codegen.
+ if (!CalleeFnDecl) {
+ if (isInOpenMPTargetExecutionDirective())
+ getASTContext().OMPTargetCalls.insert(CE);
return Call;
+ }
if (getLangOpts().OpenMP >= 50 && getLangOpts().OpenMP <= 60 &&
CalleeFnDecl->getIdentifier() &&
>From be6d4e95e9b5c7db6cc93222db08e65d7bd38cb8 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 24 Feb 2026 15:00:52 -0600
Subject: [PATCH 08/10] update sema to use new OMPTargetIndirectCall attr
---
clang/include/clang/Basic/Attr.td | 9 ++
clang/lib/CodeGen/CGExpr.cpp | 40 ++++--
clang/lib/Sema/SemaOpenMP.cpp | 20 ++-
clang/test/OpenMP/target_indirect_codegen.cpp | 58 ++++++++
.../test/api/omp_indirect_call_table_manual.c | 2 +
offload/test/api/omp_indirect_func_array.c | 124 ++++++++++++++++++
offload/test/api/omp_indirect_func_basic.c | 2 +
offload/test/api/omp_indirect_func_struct.c | 50 +++++++
offload/test/api/omp_virtual_func.cpp | 3 +
...p_virtual_func_multiple_inheritance_01.cpp | 2 +
...p_virtual_func_multiple_inheritance_02.cpp | 2 +
.../test/api/omp_virtual_func_reference.cpp | 2 +
12 files changed, 300 insertions(+), 14 deletions(-)
create mode 100644 offload/test/api/omp_indirect_func_array.c
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index b017906a8d690..7a5355d8defc0 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -4713,6 +4713,15 @@ def OMPDeclareTargetDecl : InheritableAttr {
}];
}
+def OMPTargetIndirectCall : InheritableAttr {
+ // This attribute has no spellings as it is only ever created implicitly.
+ // The attribute is added to the VarDecl of the function pointer used in
+ // an indirect call in a OMP target region.
+ let Spellings = [];
+ let SemaHandler = 0;
+ let Documentation = [InternalOnly];
+}
+
def OMPAllocateDecl : InheritableAttr {
// This attribute has no spellings as it is only ever created implicitly.
let Spellings = [];
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 8218be85484c6..5ea5dbeb60592 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -6830,18 +6830,34 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
// This is used for the indirect function case, virtual function case is
// handled in ItaniumCXXABI.cpp
if (getLangOpts().OpenMPIsTargetDevice &&
- getContext().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);
+ (!TargetDecl || !isa<FunctionDecl>(TargetDecl))) {
+ const Expr *CalleeExpr = E->getCallee()->IgnoreParenImpCasts();
+ const DeclRefExpr *DRE = nullptr;
+ while (CalleeExpr) {
+ if ((DRE = dyn_cast<DeclRefExpr>(CalleeExpr)))
+ break;
+ if (const auto *ME = dyn_cast<MemberExpr>(CalleeExpr))
+ CalleeExpr = ME->getBase()->IgnoreParenImpCasts();
+ else if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(CalleeExpr))
+ CalleeExpr = ASE->getBase()->IgnoreParenImpCasts();
+ else
+ break;
+ }
+
+ const auto *VD = DRE ? dyn_cast<VarDecl>(DRE->getDecl()) : nullptr;
+ if (VD && VD->hasAttr<OMPTargetIndirectCallAttr>()) {
+ 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;
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index b8711da4c26ea..d3f720210dd05 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7364,8 +7364,24 @@ ExprResult SemaOpenMP::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
// Mark indirect calls inside target regions, to allow for insertion of
// __llvm_omp_indirect_call_lookup calls during codegen.
if (!CalleeFnDecl) {
- if (isInOpenMPTargetExecutionDirective())
- getASTContext().OMPTargetCalls.insert(CE);
+ if (isInOpenMPTargetExecutionDirective()) {
+ Expr *E = CE->getCallee()->IgnoreParenImpCasts();
+ DeclRefExpr *DRE = nullptr;
+ while (E) {
+ if ((DRE = dyn_cast<DeclRefExpr>(E)))
+ break;
+ if (auto *ME = dyn_cast<MemberExpr>(E))
+ E = ME->getBase()->IgnoreParenImpCasts();
+ else if (auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
+ E = ASE->getBase()->IgnoreParenImpCasts();
+ else
+ break;
+ }
+ VarDecl *VD = DRE ? dyn_cast<VarDecl>(DRE->getDecl()) : nullptr;
+ if (VD && !VD->hasAttr<OMPTargetIndirectCallAttr>())
+ VD->addAttr(OMPTargetIndirectCallAttr::CreateImplicit(getASTContext()));
+ }
+
return Call;
}
diff --git a/clang/test/OpenMP/target_indirect_codegen.cpp b/clang/test/OpenMP/target_indirect_codegen.cpp
index ec249dd43b932..398d7d6407e51 100644
--- a/clang/test/OpenMP/target_indirect_codegen.cpp
+++ b/clang/test/OpenMP/target_indirect_codegen.cpp
@@ -16,6 +16,12 @@
//.
// HOST: @[[VAR:.+]] = global i8 0, align 1
+// HOST: @indirect_val = global %struct.indirect_stru { ptr @_Z3bazv }, align 8
+// HOST: @indirect_nested_val = global %struct.indirect_stru_nested { %struct.indirect_stru { ptr @_Z3bazv } }, align 8
+// HOST: @indirect_baz = global ptr @_Z3bazv, align 8
+// HOST: @indirect_bar = global ptr @_ZL3barv, align 8
+// HOST: @indirect_foo = global ptr @_Z3foov, align 8
+// HOST: @indirect_array = global [3 x ptr] [ptr @_Z3foov, ptr @_ZL3barv, ptr @_Z3bazv], align 8
// HOST: @[[FOO_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[FOO_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_foo_l[0-9]+]]\00"
// HOST: @.offloading.entry.[[FOO_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3foov, ptr @[[FOO_ENTRY_NAME]], i64 8, i64 0, ptr null }
// HOST: @[[BAZ_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[BAZ_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_baz_l[0-9]+]]\00"
@@ -47,8 +53,60 @@ void disabled() { };
char var = 0;
#pragma omp declare target to(var) indirect
+struct indirect_stru {
+ void (*arg)();
+};
+struct indirect_stru_nested {
+ struct indirect_stru nested;
+};
+
+struct indirect_stru indirect_val = { .arg = baz };
+struct indirect_stru_nested indirect_nested_val = { .nested = { .arg = baz } };
+
+void (*indirect_baz)() = baz;
+void (*indirect_bar)() = bar;
+void (*indirect_foo)() = foo;
+void (*indirect_array[3])() = { foo, bar, baz };
+
+
+int main() {
+ #pragma omp target map(indirect_baz,indirect_bar,indirect_foo,var,indirect_val,indirect_val.arg, indirect_array, indirect_array[0:3], indirect_nested_val, indirect_nested_val.nested.arg)
+ {
+ indirect_foo();
+ indirect_bar();
+ indirect_baz();
+ indirect_val.arg();
+ indirect_nested_val.nested.arg();
+ indirect_array[0]();
+ indirect_array[1]();
+ indirect_array[2]();
+ }
+}
#endif
+// DEVICE-LABEL: define {{.*}}void @__omp_offloading_{{.+}}_main_l{{[0-9]+}}(
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: getelementptr inbounds nuw %struct.indirect_stru,
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: getelementptr inbounds nuw %struct.indirect_stru_nested,
+// DEVICE: getelementptr inbounds nuw %struct.indirect_stru,
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: getelementptr inbounds {{.*}}[3 x ptr{{[^]]*}}],
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: getelementptr inbounds {{.*}}[3 x ptr{{[^]]*}}],
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
+// DEVICE: getelementptr inbounds {{.*}}[3 x ptr{{[^]]*}}],
+// DEVICE: call {{.*}}@__llvm_omp_indirect_call_lookup(
+// DEVICE: call {{.*}}void %{{.+}}()
//.
// HOST-DAG: !{{[0-9]+}} = !{i32 1, !"[[FOO_NAME]]", i32 8, i32 0}
// HOST-DAG: !{{[0-9]+}} = !{i32 1, !"[[BAZ_NAME]]", i32 8, i32 1}
diff --git a/offload/test/api/omp_indirect_call_table_manual.c b/offload/test/api/omp_indirect_call_table_manual.c
index e958d47d69dad..550c9fc4d201c 100644
--- a/offload/test/api/omp_indirect_call_table_manual.c
+++ b/offload/test/api/omp_indirect_call_table_manual.c
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// XFAIL: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/api/omp_indirect_func_array.c b/offload/test/api/omp_indirect_func_array.c
new file mode 100644
index 0000000000000..1064051d412fd
--- /dev/null
+++ b/offload/test/api/omp_indirect_func_array.c
@@ -0,0 +1,124 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
+
+#include <assert.h>
+#include <stdio.h>
+
+#define TEST_VAL 5
+
+#pragma omp declare target indirect
+int func_a(int x) { return x + 1; }
+int func_b(int x) { return x + 2; }
+int func_c(int x) { return x + 3; }
+int func_d(int x) { return x * 2; }
+int func_e(int x) { return x * 3; }
+int func_f(int x) { return x * 4; }
+#pragma omp end declare target
+
+void test_array_explicit_mapping() {
+ int (*local_fptr_array[3])(int) = {func_a, func_b, func_c};
+
+ int results[3];
+ int expected[3];
+
+ expected[0] = func_a(TEST_VAL);
+ expected[1] = func_b(TEST_VAL);
+ expected[2] = func_c(TEST_VAL);
+
+#pragma omp target map(local_fptr_array, local_fptr_array[0 : 3]) \
+ map(from : results)
+ {
+ for (int i = 0; i < 3; i++) {
+ results[i] = local_fptr_array[i](TEST_VAL);
+ }
+ }
+
+ for (int i = 0; i < 3; i++) {
+ assert(results[i] == expected[i] &&
+ "Error: local array function pointer returned incorrect value on "
+ "device");
+ }
+
+ // Change function pointers and re-test
+ local_fptr_array[0] = func_d;
+ local_fptr_array[1] = func_e;
+ local_fptr_array[2] = func_f;
+
+ expected[0] = func_d(TEST_VAL);
+ expected[1] = func_e(TEST_VAL);
+ expected[2] = func_f(TEST_VAL);
+
+#pragma omp target map(local_fptr_array, local_fptr_array[0 : 3]) \
+ map(from : results)
+ {
+ for (int i = 0; i < 3; i++) {
+ results[i] = local_fptr_array[i](TEST_VAL);
+ }
+ }
+
+ for (int i = 0; i < 3; i++) {
+ assert(results[i] == expected[i] &&
+ "Error: local array function pointer returned incorrect value on "
+ "device after update");
+ }
+}
+
+struct with_fptr_array {
+ int buffer;
+ int (*fptrs[3])(int);
+};
+
+void test_struct_containing_array() {
+ struct with_fptr_array val = {.buffer = 0, .fptrs = {func_a, func_b, func_c}};
+
+ int results[3];
+ int expected[3];
+
+ expected[0] = func_a(TEST_VAL);
+ expected[1] = func_b(TEST_VAL);
+ expected[2] = func_c(TEST_VAL);
+
+#pragma omp target map(val, val.fptrs[0 : 3]) map(from : results)
+ {
+ results[0] = val.fptrs[0](TEST_VAL);
+ results[1] = val.fptrs[1](TEST_VAL);
+ results[2] = val.fptrs[2](TEST_VAL);
+ }
+
+ for (int i = 0; i < 3; i++) {
+ assert(results[i] == expected[i] &&
+ "Error: struct array function pointer returned incorrect value");
+ }
+
+ // Update and re-test
+ val.fptrs[0] = func_d;
+ val.fptrs[1] = func_e;
+ val.fptrs[2] = func_f;
+
+ expected[0] = func_d(TEST_VAL);
+ expected[1] = func_e(TEST_VAL);
+ expected[2] = func_f(TEST_VAL);
+
+#pragma omp target map(val, val.fptrs[0 : 3]) map(from : results)
+ {
+ results[0] = val.fptrs[0](TEST_VAL);
+ results[1] = val.fptrs[1](TEST_VAL);
+ results[2] = val.fptrs[2](TEST_VAL);
+ }
+
+ for (int i = 0; i < 3; i++) {
+ assert(results[i] == expected[i] &&
+ "Error: struct array function pointer returned incorrect value "
+ "after update");
+ }
+}
+
+int main() {
+ test_array_explicit_mapping();
+ test_struct_containing_array();
+
+ // CHECK: PASS
+ printf("PASS\n");
+ return 0;
+}
diff --git a/offload/test/api/omp_indirect_func_basic.c b/offload/test/api/omp_indirect_func_basic.c
index aaf0ccf371015..b2e29b92f13b3 100644
--- a/offload/test/api/omp_indirect_func_basic.c
+++ b/offload/test/api/omp_indirect_func_basic.c
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
diff --git a/offload/test/api/omp_indirect_func_struct.c b/offload/test/api/omp_indirect_func_struct.c
index f9e1489e4b4ea..527c9e0b48c5f 100644
--- a/offload/test/api/omp_indirect_func_struct.c
+++ b/offload/test/api/omp_indirect_func_struct.c
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
@@ -27,10 +29,57 @@ typedef struct {
#pragma omp declare mapper(indirect_stru_mapped s) \
map(s, s.indirect0_ptr, s.indirect1_ptr)
+struct nested_stru {
+ struct indirect_stru inner;
+};
+
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};
+struct nested_stru global_nested_val = {
+ .inner = {.buffer = 0,
+ .indirect0 = indirect_base_arg,
+ .indirect1 = indirect_base}};
+
+void test_nested_struct() {
+ int indirect0_ret, indirect1_ret;
+ int indirect0_base = indirect_base_arg(TEST_VAL);
+ int indirect1_base = indirect_base();
+
+#pragma omp target map(global_nested_val, global_nested_val.inner.indirect0, \
+ global_nested_val.inner.indirect1) \
+ map(from : indirect0_ret, indirect1_ret)
+ {
+ indirect0_ret = global_nested_val.inner.indirect0(TEST_VAL);
+ indirect1_ret = global_nested_val.inner.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base &&
+ "Error: indirect0 returned incorrect value on device");
+ assert(indirect1_ret == indirect1_base &&
+ "Error: indirect1 returned incorrect value on device");
+
+ // Change function pointers and re-test
+ global_nested_val.inner.indirect0 = direct_arg;
+ global_nested_val.inner.indirect1 = direct;
+
+ indirect0_base = direct_arg(TEST_VAL);
+ indirect1_base = direct();
+
+#pragma omp target map(global_nested_val, global_nested_val.inner.indirect0, \
+ global_nested_val.inner.indirect1) \
+ map(from : indirect0_ret, indirect1_ret)
+ {
+ indirect0_ret = global_nested_val.inner.indirect0(TEST_VAL);
+ indirect1_ret = global_nested_val.inner.indirect1();
+ }
+
+ assert(indirect0_ret == indirect0_base &&
+ "Error: indirect0 returned incorrect value on device after update");
+ assert(indirect1_ret == indirect1_base &&
+ "Error: indirect1 returned incorrect value on device after update");
+}
void test_global_struct_explicit_mapping() {
int indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
@@ -261,6 +310,7 @@ void test_local_struct_user_mapper() {
}
int main() {
+ test_nested_struct();
test_global_struct_explicit_mapping();
test_local_struct_explicit_mapping();
test_global_struct_user_mapper();
diff --git a/offload/test/api/omp_virtual_func.cpp b/offload/test/api/omp_virtual_func.cpp
index 70d35e07cfbb3..714ba33a513df 100644
--- a/offload/test/api/omp_virtual_func.cpp
+++ b/offload/test/api/omp_virtual_func.cpp
@@ -1,4 +1,7 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
+
#include <assert.h>
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
index 307ee78d01202..373d6fd2af06b 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
index 1aba0304d17e1..ae61bfc02d060 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
diff --git a/offload/test/api/omp_virtual_func_reference.cpp b/offload/test/api/omp_virtual_func_reference.cpp
index 7fb25dbd5c626..deb54b291438c 100644
--- a/offload/test/api/omp_virtual_func_reference.cpp
+++ b/offload/test/api/omp_virtual_func_reference.cpp
@@ -1,4 +1,6 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+// REQUIRES: gpu
#include <assert.h>
#include <omp.h>
>From bb9e61a5b72a44bf38f713f1d836cfa38df9c98b Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Wed, 25 Feb 2026 17:00:01 -0600
Subject: [PATCH 09/10] add ASTMutationListener
---
clang/include/clang/AST/ASTMutationListener.h | 5 +++++
clang/include/clang/Serialization/ASTWriter.h | 1 +
clang/lib/Frontend/MultiplexConsumer.cpp | 6 ++++++
clang/lib/Sema/SemaOpenMP.cpp | 5 ++++-
clang/lib/Serialization/ASTCommon.h | 1 +
clang/lib/Serialization/ASTReaderDecl.cpp | 5 +++++
clang/lib/Serialization/ASTWriter.cpp | 15 +++++++++++++++
7 files changed, 37 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/ASTMutationListener.h b/clang/include/clang/AST/ASTMutationListener.h
index c8448a25c23a4..947791f4a506c 100644
--- a/clang/include/clang/AST/ASTMutationListener.h
+++ b/clang/include/clang/AST/ASTMutationListener.h
@@ -153,6 +153,11 @@ class ASTMutationListener {
/// \param D the declaration marked as a variable with OpenMP allocator.
virtual void DeclarationMarkedOpenMPAllocate(const Decl *D, const Attr *A) {}
+ /// A declaration is marked as an OpenMP indirect call target.
+ ///
+ /// \param D the declaration marked as an indirect call target.
+ virtual void DeclarationMarkedOpenMPIndirectCall(const Decl *D) {}
+
/// A definition has been made visible by being redefined locally.
///
/// \param D The definition that was previously not visible.
diff --git a/clang/include/clang/Serialization/ASTWriter.h b/clang/include/clang/Serialization/ASTWriter.h
index d3029373ed2f7..a00a416c260e5 100644
--- a/clang/include/clang/Serialization/ASTWriter.h
+++ b/clang/include/clang/Serialization/ASTWriter.h
@@ -973,6 +973,7 @@ class ASTWriter : public ASTDeserializationListener,
void DeclarationMarkedOpenMPDeclareTarget(const Decl *D,
const Attr *Attr) override;
void DeclarationMarkedOpenMPAllocate(const Decl *D, const Attr *A) override;
+ void DeclarationMarkedOpenMPIndirectCall(const Decl *D) override;
void RedefinedHiddenDefinition(const NamedDecl *D, Module *M) override;
void AddedAttributeToRecord(const Attr *Attr,
const RecordDecl *Record) override;
diff --git a/clang/lib/Frontend/MultiplexConsumer.cpp b/clang/lib/Frontend/MultiplexConsumer.cpp
index f5f8848798a35..abbd5eef7c16b 100644
--- a/clang/lib/Frontend/MultiplexConsumer.cpp
+++ b/clang/lib/Frontend/MultiplexConsumer.cpp
@@ -120,6 +120,7 @@ class MultiplexASTMutationListener : public ASTMutationListener {
void DeclarationMarkedUsed(const Decl *D) override;
void DeclarationMarkedOpenMPThreadPrivate(const Decl *D) override;
void DeclarationMarkedOpenMPAllocate(const Decl *D, const Attr *A) override;
+ void DeclarationMarkedOpenMPIndirectCall(const Decl *D) override;
void DeclarationMarkedOpenMPDeclareTarget(const Decl *D,
const Attr *Attr) override;
void RedefinedHiddenDefinition(const NamedDecl *D, Module *M) override;
@@ -240,6 +241,11 @@ void MultiplexASTMutationListener::DeclarationMarkedOpenMPAllocate(
for (ASTMutationListener *L : Listeners)
L->DeclarationMarkedOpenMPAllocate(D, A);
}
+void MultiplexASTMutationListener::DeclarationMarkedOpenMPIndirectCall(
+ const Decl *D) {
+ for (ASTMutationListener *L : Listeners)
+ L->DeclarationMarkedOpenMPIndirectCall(D);
+}
void MultiplexASTMutationListener::DeclarationMarkedOpenMPDeclareTarget(
const Decl *D, const Attr *Attr) {
for (auto *L : Listeners)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index d3f720210dd05..d38d92c32419d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7378,8 +7378,11 @@ ExprResult SemaOpenMP::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
break;
}
VarDecl *VD = DRE ? dyn_cast<VarDecl>(DRE->getDecl()) : nullptr;
- if (VD && !VD->hasAttr<OMPTargetIndirectCallAttr>())
+ if (VD && !VD->hasAttr<OMPTargetIndirectCallAttr>()) {
VD->addAttr(OMPTargetIndirectCallAttr::CreateImplicit(getASTContext()));
+ if (ASTMutationListener *ML = getASTContext().getASTMutationListener())
+ ML->DeclarationMarkedOpenMPIndirectCall(VD);
+ }
}
return Call;
diff --git a/clang/lib/Serialization/ASTCommon.h b/clang/lib/Serialization/ASTCommon.h
index 23d3954f257e7..e2835168cbf9e 100644
--- a/clang/lib/Serialization/ASTCommon.h
+++ b/clang/lib/Serialization/ASTCommon.h
@@ -39,6 +39,7 @@ enum class DeclUpdateKind {
StaticLocalNumber,
DeclMarkedOpenMPThreadPrivate,
DeclMarkedOpenMPAllocate,
+ DeclMarkedOpenMPIndirectCall,
DeclMarkedOpenMPDeclareTarget,
DeclExported,
AddedAttrToRecord,
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index f8e9caa3f5d1d..f33915a7ee1d6 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -4973,6 +4973,11 @@ void ASTDeclReader::UpdateDecl(Decl *D) {
break;
}
+ case DeclUpdateKind::DeclMarkedOpenMPIndirectCall:
+ D->addAttr(OMPTargetIndirectCallAttr::CreateImplicit(Reader.getContext(),
+ readSourceRange()));
+ break;
+
case DeclUpdateKind::DeclExported: {
unsigned SubmoduleID = readSubmoduleID();
auto *Exported = cast<NamedDecl>(D);
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index f9176b7e68f73..5829e7beace50 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -6611,6 +6611,11 @@ void ASTWriter::WriteDeclUpdatesBlocks(ASTContext &Context,
break;
}
+ case DeclUpdateKind::DeclMarkedOpenMPIndirectCall:
+ Record.AddSourceRange(
+ D->getAttr<OMPTargetIndirectCallAttr>()->getRange());
+ break;
+
case DeclUpdateKind::DeclMarkedOpenMPDeclareTarget:
Record.push_back(D->getAttr<OMPDeclareTargetDeclAttr>()->getMapType());
Record.AddSourceRange(
@@ -7799,6 +7804,16 @@ void ASTWriter::DeclarationMarkedOpenMPAllocate(const Decl *D, const Attr *A) {
DeclUpdate(DeclUpdateKind::DeclMarkedOpenMPAllocate, A));
}
+void ASTWriter::DeclarationMarkedOpenMPIndirectCall(const Decl *D) {
+ if (Chain && Chain->isProcessingUpdateRecords()) return;
+ assert(!WritingAST && "Already writing the AST!");
+ if (!D->isFromASTFile())
+ return;
+
+ DeclUpdates[D].push_back(
+ DeclUpdate(DeclUpdateKind::DeclMarkedOpenMPIndirectCall));
+}
+
void ASTWriter::DeclarationMarkedOpenMPDeclareTarget(const Decl *D,
const Attr *Attr) {
if (Chain && Chain->isProcessingUpdateRecords()) return;
>From 623545e0db7f5d0a535c9585740f8b5357c9bfe0 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Mon, 2 Mar 2026 09:34:17 -0600
Subject: [PATCH 10/10] update on feedback
---
clang/include/clang/AST/ASTContext.h | 4 ----
1 file changed, 4 deletions(-)
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index dc27e2bd5d1d4..68205dd1c1fd9 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -1370,10 +1370,6 @@ class ASTContext : public RefCountedBase<ASTContext> {
/// are stored here.
llvm::DenseMap<const CXXMethodDecl *, CXXCastPath> LambdaCastPaths;
- /// Keep track of indirect call expressions within OpenMP target regions.
- /// Used to instert calls to __llvm_omp_indirect_call_lookup during codegen.
- llvm::DenseSet<const CallExpr *> OMPTargetCalls;
-
ASTContext(LangOptions &LOpts, SourceManager &SM, IdentifierTable &idents,
SelectorTable &sels, Builtin::Context &builtins,
TranslationUnitKind TUKind);
More information about the cfe-commits
mailing list