Author: Jason-VanBeusekom Date: 2026-01-05T13:25:03-05:00 New Revision: 7a174c91f3c7630ad5cd75622846a59dee2fc1b9
URL: https://github.com/llvm/llvm-project/commit/7a174c91f3c7630ad5cd75622846a59dee2fc1b9 DIFF: https://github.com/llvm/llvm-project/commit/7a174c91f3c7630ad5cd75622846a59dee2fc1b9.diff LOG: [OpenMP][clang] Register Vtables on device for indirect calls (#159856) Added: clang/test/OpenMP/target_vtable_codegen_container.cpp clang/test/OpenMP/target_vtable_codegen_explicit.cpp clang/test/OpenMP/target_vtable_codegen_implicit_namespace.cpp clang/test/OpenMP/target_vtable_codegen_memberexpr_codegen.cpp clang/test/OpenMP/target_vtable_codegen_mult_inherritence.cpp clang/test/OpenMP/target_vtable_codegen_nested.cpp Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/lib/CodeGen/CGOpenMPRuntime.h clang/lib/CodeGen/CGStmtOpenMP.cpp clang/lib/CodeGen/CGVTables.cpp clang/lib/CodeGen/CGVTables.h llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index b8ee701c482bb..01661ad54ee2f 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -1776,12 +1776,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; + 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); +} + +void CGOpenMPRuntime::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.contains(CXXRecord)) { + auto Res = CGM.getOpenMPRuntime().VTableDeclMap.try_emplace(CXXRecord, VD); + if (Res.second) { + CGM.EmitVTable(CXXRecord); + CodeGenVTables VTables = CGM.getVTables(); + llvm::GlobalVariable *VTablesAddr = VTables.GetAddrOfVTable(CXXRecord); + assert(VTablesAddr && "Expected non-null VTable address"); + 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); + } + } + } +} + +void CGOpenMPRuntime::registerVTable(const OMPExecutableDirective &D) { + // Register VTable by scanning through the map clause of OpenMP target region. + // 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()); + } else if (auto *MRE = dyn_cast<MemberExpr>(E)) { + if (auto *BaseDRE = dyn_cast<DeclRefExpr>(MRE->getBase())) { + if (auto *BaseVD = dyn_cast<VarDecl>(BaseDRE->getDecl())) + VD = BaseVD; + } + } + return std::pair<CXXRecordDecl *, const VarDecl *>(CXXRecord, VD); + }; + // Collect VTable from OpenMP map clause. + for (const auto *C : D.getClausesOfKind<OMPMapClause>()) { + for (const auto *E : C->varlist()) { + auto DeclPair = GetVTableDecl(E); + // Ensure VD is not null + if (DeclPair.second) + emitAndRegisterVTable(CGM, DeclPair.first, DeclPair.second); + } + } +} + Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF, QualType VarType, StringRef Name) { @@ -2845,6 +2959,13 @@ void CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata() { case llvm::OpenMPIRBuilder::EMIT_MD_GLOBAL_VAR_LINK_ERROR: { CGM.getDiags().Report(diag::err_target_var_offloading_entry_incorrect); } break; + case llvm::OpenMPIRBuilder::EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR: { + unsigned DiagID = CGM.getDiags().getCustomDiagID( + DiagnosticsEngine::Error, "Offloading entry for indirect declare " + "target variable is incorrect: the " + "address is invalid."); + CGM.getDiags().Report(DiagID); + } break; } }; @@ -6249,6 +6370,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper( CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr); } } + registerVTable(D); } /// Checks if the expression is constant or does not have non-trivial function @@ -10903,6 +11025,17 @@ 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 (auto *E = dyn_cast<OMPExecutableDirective>(S); + E && isOpenMPTargetDataManagementDirective(E->getDirectiveKind())) { + // 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 6bfd7d6a590b9..a81d3830a8035 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::SmallDenseMap<CXXRecordDecl *, const VarDecl *> VTableDeclMap; + public: explicit CGOpenMPRuntime(CodeGenModule &CGM); virtual ~CGOpenMPRuntime() {} @@ -1111,6 +1114,23 @@ 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); + + /// Emit and register VTable for the C++ class in OpenMP offload entry. + /// \param CXXRecord C++ class decl. + /// \param VD Variable decl which holds VTable. + virtual void emitAndRegisterVTable(CodeGenModule &CGM, + CXXRecordDecl *CXXRecord, + const VarDecl *VD); + /// 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 fee275c018294..fa1e7361adbff 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -7659,6 +7659,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 f0d7e7003f2d9..46f00192c5735 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/test/OpenMP/target_vtable_codegen_container.cpp b/clang/test/OpenMP/target_vtable_codegen_container.cpp new file mode 100644 index 0000000000000..9fd4c6b736163 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_container.cpp @@ -0,0 +1,42 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + +// CHECK-DAG: @_ZTV7Derived +// CHECK-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; +} diff --git a/clang/test/OpenMP/target_vtable_codegen_explicit.cpp b/clang/test/OpenMP/target_vtable_codegen_explicit.cpp new file mode 100644 index 0000000000000..001ed8fdd9cd7 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_explicit.cpp @@ -0,0 +1,48 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + +// Make sure both host and device compilation emit vtable for Dervied +// CHECK-DAG: $_ZN7DerivedD1Ev = comdat any +// CHECK-DAG: $_ZN7DerivedD0Ev = comdat any +// CHECK-DAG: $_ZN7Derived5BaseAEi = comdat any +// CHECK-DAG: $_ZN7Derived8DerivedBEv = comdat any +// CHECK-DAG: $_ZN7DerivedD2Ev = comdat any +// CHECK-DAG: $_ZN4BaseD2Ev = comdat any +// CHECK-DAG: $_ZTV7Derived = comdat any +class Base { +public: + + virtual ~Base() = default; + + virtual void BaseA(int a) { } +}; + +// CHECK: @_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; +} diff --git a/clang/test/OpenMP/target_vtable_codegen_implicit_namespace.cpp b/clang/test/OpenMP/target_vtable_codegen_implicit_namespace.cpp new file mode 100644 index 0000000000000..364c55cd07985 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_implicit_namespace.cpp @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + +namespace { + +// Make sure both host and device compilation emit vtable for Dervied +// CHECK-DAG: @_ZTVN12_GLOBAL__N_17DerivedE +// CHECK-DAG: @_ZN12_GLOBAL__N_17DerivedD1Ev +// CHECK-DAG: @_ZN12_GLOBAL__N_17DerivedD0Ev +// CHECK-DAG: @_ZN12_GLOBAL__N_17Derived5BaseAEi +// CHECK-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; +} diff --git a/clang/test/OpenMP/target_vtable_codegen_memberexpr_codegen.cpp b/clang/test/OpenMP/target_vtable_codegen_memberexpr_codegen.cpp new file mode 100644 index 0000000000000..0535ba1dec741 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_memberexpr_codegen.cpp @@ -0,0 +1,56 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + + +// CHECK-DAG: $_ZN4Base5BaseAEi = comdat any +// CHECK-DAG: $_ZN7Derived5BaseAEi = comdat any +// CHECK-DAG: $_ZN7Derived8DerivedBEv = comdat any +// CHECK-DAG: $_ZN4BaseD1Ev = comdat any +// CHECK-DAG: $_ZN4BaseD0Ev = comdat any +// CHECK-DAG: $_ZN7DerivedD1Ev = comdat any +// CHECK-DAG: $_ZN7DerivedD0Ev = comdat any +// CHECK-DAG: $_ZN4BaseD2Ev = comdat any +// CHECK-DAG: $_ZN7DerivedD2Ev = comdat any +// CHECK-DAG: $_ZTV4Base = comdat any +// CHECK-DAG: $_ZTV7Derived = comdat any +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; +}; + +struct VirtualContainer { + Base baseObj; + Derived derivedObj; + Base *basePtr; +}; + +int main() { + VirtualContainer container; + container.basePtr = &container.derivedObj; + int a = 50; +#pragma omp target map(container.baseObj, container.derivedObj, \ + container.basePtr[ : 1]) + { + container.baseObj.BaseA(a); + container.derivedObj.BaseA(a); + container.derivedObj.DerivedB(); + container.basePtr->BaseA(a); + } + return 0; +} diff --git a/clang/test/OpenMP/target_vtable_codegen_mult_inherritence.cpp b/clang/test/OpenMP/target_vtable_codegen_mult_inherritence.cpp new file mode 100644 index 0000000000000..3069a4994a479 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_mult_inherritence.cpp @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + +// CHECK-DAG: @_ZTV6Base_1 +// CHECK-DAG: @_ZTV7Derived +// CHECK-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; +} diff --git a/clang/test/OpenMP/target_vtable_codegen_nested.cpp b/clang/test/OpenMP/target_vtable_codegen_nested.cpp new file mode 100644 index 0000000000000..1ece83d60ac58 --- /dev/null +++ b/clang/test/OpenMP/target_vtable_codegen_nested.cpp @@ -0,0 +1,82 @@ +// RUN: %clang_cc1 -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 -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 +// expected-no-diagnostics + +// CHECK-DAG: @_ZTV3Car +// CHECK-DAG: @_ZTV6Engine +// CHECK-DAG: @_ZTV6Wheels +// CHECK-DAG: @_ZTV7Vehicle +// CHECK-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; +} diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 05d8a7dd168a3..c8db40d3cf51b 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -396,6 +396,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 @@ -2761,7 +2763,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 5e4d4c7e49776..b6a3d9e66fb9c 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -10634,6 +10634,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; } @@ -10643,12 +10650,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 @@ -11085,7 +11097,9 @@ 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()); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
