https://github.com/Jason-VanBeusekom updated https://github.com/llvm/llvm-project/pull/159857
>From f783b0889dc94fda28ac9086899332015c441db4 Mon Sep 17 00:00:00 2001 From: "[email protected]" <[email protected]> Date: Fri, 12 Sep 2025 14:08:24 -0500 Subject: [PATCH 1/3] [OpenMP][clang] Indirect and Virtual function call mapping from host to device This patch implements the CodeGen logic for calling __llvm_omp_indirect_call_lookup on the device when an indirect function call or a virtual function call is made within an OpenMP target region. --------- Co-authored-by: Youngsuk Kim --- clang/lib/CodeGen/CGExpr.cpp | 20 + clang/lib/CodeGen/CGOpenMPRuntime.cpp | 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 <[email protected]> Date: Wed, 1 Oct 2025 13:18:01 -0500 Subject: [PATCH 2/3] Updates based on feedback --- clang/lib/CodeGen/CGExpr.cpp | 7 +++---- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 3 +-- clang/lib/CodeGen/ItaniumCXXABI.cpp | 11 +++++------ .../target_vtable_omp_indirect_call_lookup.cpp | 8 ++++---- offload/test/api/omp_indirect_call.c | 12 ++++++------ openmp/device/src/Misc.cpp | 2 +- 6 files changed, 20 insertions(+), 23 deletions(-) diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index 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 <[email protected]> Date: Tue, 6 Jan 2026 11:46:30 -0600 Subject: [PATCH 3/3] 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(); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
