[clang] 7a174c9 - [OpenMP][clang] Register Vtables on device for indirect calls (#159856)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 5 10:25:08 PST 2026
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());
More information about the cfe-commits
mailing list