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