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/Open...
[truncated]

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index bf91da314..8b23d0470 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1832,36 +1832,31 @@ void CGOpenMPRuntime::registerVTableOffloadEntry(llvm::GlobalVariable *VTable,
 }
 
 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.find(CXXRecord) ==
-                CGM.getOpenMPRuntime().VTableDeclMap.end()) {
-          CGM.getOpenMPRuntime().VTableDeclMap.try_emplace(CXXRecord, VD);
-          CGM.EmitVTable(CXXRecord);
-          CodeGenVTables VTables = CGM.getVTables();
-          llvm::GlobalVariable *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);
-
-          }
-        }
-      };
-
+                                            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);
+    CodeGenVTables VTables = CGM.getVTables();
+    llvm::GlobalVariable *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);
+    }
+  }
+};
 
 void CGOpenMPRuntime::registerVTable(const OMPExecutableDirective &D) {
   // Register VTable by scanning through the map clause of OpenMP target region.
@@ -10473,15 +10468,15 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
   // 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(
-            dyn_cast<OMPExecutableDirective>(S)->getDirectiveKind())) {
-      auto &E = *dyn_cast<OMPExecutableDirective>(S);
-      // Don't need to check if it's device compile
-      // since scanForTargetRegionsFunctions currently only called
-      // in device compilation.
-      registerVTable(E);
-    }
+  if (isa<OMPExecutableDirective>(S) &&
+      isOpenMPTargetDataManagementDirective(
+          dyn_cast<OMPExecutableDirective>(S)->getDirectiveKind())) {
+    auto &E = *dyn_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 =
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 4ace1abcb..49dcba4b7 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -754,7 +754,7 @@ public:
   // i32 @__isPlatformVersionAtLeast(i32, i32, i32, i32)
   llvm::FunctionCallee IsPlatformVersionAtLeastFn = nullptr;
 
-  //  Store indirect CallExprs that are within an omp target region 
+  //  Store indirect CallExprs that are within an omp target region
   llvm::SmallPtrSet<const CallExpr *, 16> OMPTargetCalls;
 
   InstrProfStats &getPGOStats() { return PGOStats; }
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 563545060..1a8c52894 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10260,8 +10260,8 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
       case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect:
       case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable:
         if (!CE->getAddress()) {
-            ErrorFn(EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR, E.second);
-            continue;
+          ErrorFn(EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR, E.second);
+          continue;
         }
         break;
       default:
@@ -10715,7 +10715,8 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
       return;
     }
     if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect ||
-        Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable)
+        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 8dfda40ca..613e73b04 100644
--- a/offload/include/omptarget.h
+++ b/offload/include/omptarget.h
@@ -94,7 +94,7 @@ 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.
+  /// Mark the entry global as being an indirect vtable.
   OMP_DECLARE_TARGET_INDIRECT_VTABLE = 0x20,
 };
 
diff --git a/offload/test/api/omp_indirect_call_table_manual.c b/offload/test/api/omp_indirect_call_table_manual.c
index 9c6fd4ca8..e958d47d6 100644
--- a/offload/test/api/omp_indirect_call_table_manual.c
+++ b/offload/test/api/omp_indirect_call_table_manual.c
@@ -89,7 +89,7 @@ int main() {
   myClass obj_baz = {dispatchTable + 2};
   int aaa = 0;
 
-#pragma omp target map(aaa) map (to: obj_foo, obj_bar, obj_baz)
+#pragma omp target map(aaa) map(to : obj_foo, obj_bar, obj_baz)
   {
     // Lookup
     fptr_t *foo_ptr = __llvm_omp_indirect_call_lookup(obj_foo.dispatchPtr);