Index: clang/include/clang/Sema/ScopeInfo.h
===================================================================
--- clang/include/clang/Sema/ScopeInfo.h
+++ clang/include/clang/Sema/ScopeInfo.h
@@ -756,13 +756,15 @@
   unsigned short CapRegionKind;
 
   unsigned short OpenMPLevel;
+  unsigned short OpenMPCaptureLevel;
 
   CapturedRegionScopeInfo(DiagnosticsEngine &Diag, Scope *S, CapturedDecl *CD,
                           RecordDecl *RD, ImplicitParamDecl *Context,
                           CapturedRegionKind K, unsigned OpenMPLevel)
       : CapturingScopeInfo(Diag, ImpCap_CapturedRegion),
         TheCapturedDecl(CD), TheRecordDecl(RD), TheScope(S),
-        ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel) {
+        ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel),
+        OpenMPCaptureLevel(0) {
     Kind = SK_CapturedRegion;
   }
 
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -8980,10 +8980,6 @@
   /// Returns OpenMP nesting level for current directive.
   unsigned getOpenMPNestingLevel() const;
 
-  /// Adjusts the function scopes index for the target-based regions.
-  void adjustOpenMPTargetScopeIndex(unsigned &FunctionScopesIndex,
-                                    unsigned Level) const;
-
   /// Push new OpenMP function region for non-capturing function.
   void pushOpenMPFunctionRegion();
 
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -2108,10 +2108,16 @@
 
 void Sema::PushCapturedRegionScope(Scope *S, CapturedDecl *CD, RecordDecl *RD,
                                    CapturedRegionKind K) {
-  CapturingScopeInfo *CSI = new CapturedRegionScopeInfo(
+  CapturedRegionScopeInfo *CSI = new CapturedRegionScopeInfo(
       getDiagnostics(), S, CD, RD, CD->getContextParam(), K,
       (getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0);
   CSI->ReturnType = Context.VoidTy;
+  if (getLangOpts().OpenMP && K == CR_OpenMP) {
+    if (auto *P = dyn_cast<CapturedRegionScopeInfo>(FunctionScopes.back())) {
+      if (P->CapRegionKind == CR_OpenMP && CSI->OpenMPLevel == P->OpenMPLevel)
+        CSI->OpenMPCaptureLevel = P->OpenMPCaptureLevel + 1;
+    }
+  }
   FunctionScopes.push_back(CSI);
 }
 
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -15784,7 +15784,7 @@
           // target region, therefore we need to propagate the capture from the
           // enclosing region. Therefore, the capture is not initially nested.
           if (IsTargetCap)
-            adjustOpenMPTargetScopeIndex(FunctionScopesIndex, RSI->OpenMPLevel);
+            FunctionScopesIndex -= RSI->OpenMPCaptureLevel + 1;
 
           if (IsTargetCap || IsOpenMPPrivateDecl) {
             Nested = !IsTargetCap;
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -1853,13 +1853,6 @@
   return nullptr;
 }
 
-void Sema::adjustOpenMPTargetScopeIndex(unsigned &FunctionScopesIndex,
-                                        unsigned Level) const {
-  SmallVector<OpenMPDirectiveKind, 4> Regions;
-  getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level));
-  FunctionScopesIndex -= Regions.size();
-}
-
 void Sema::startOpenMPLoop() {
   assert(LangOpts.OpenMP && "OpenMP must be enabled.");
   if (isOpenMPLoopDirective(DSAStack->getCurrentDirective()))
@@ -3408,6 +3401,7 @@
   OMPScheduleClause *SC = nullptr;
   SmallVector<const OMPLinearClause *, 4> LCs;
   SmallVector<const OMPClauseWithPreInit *, 4> PICs;
+  SmallVector<OMPMapClause *, 4> MCs;
   // This is required for proper codegen.
   for (OMPClause *Clause : Clauses) {
     if (isOpenMPTaskingDirective(DSAStack->getCurrentDirective()) &&
@@ -3447,6 +3441,8 @@
       OC = cast<OMPOrderedClause>(Clause);
     else if (Clause->getClauseKind() == OMPC_linear)
       LCs.push_back(cast<OMPLinearClause>(Clause));
+    else if (Clause->getClauseKind() == OMPC_map)
+      MCs.push_back(cast<OMPMapClause>(Clause));
   }
   // OpenMP, 2.7.1 Loop Construct, Restrictions
   // The nonmonotonic modifier cannot be specified if an ordered clause is
@@ -3503,6 +3499,14 @@
         }
       }
     }
+    if (ThisCaptureRegion == OMPD_target) {
+      for (OMPMapClause *MC : MCs) {
+        for (ValueDecl *D : MC->all_decls()) {
+          if (auto *VD = dyn_cast_or_null<VarDecl>(D))
+            MarkVariableReferenced(VD->getLocation(), VD);
+        }
+      }
+    }
     if (++CompletedRegions == CaptureRegions.size())
       DSAStack->setBodyComplete();
     SR = ActOnCapturedRegionEnd(SR.get());
Index: clang/test/OpenMP/target_map_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen.cpp
+++ clang/test/OpenMP/target_map_codegen.cpp
@@ -5329,5 +5329,125 @@
 // CK31: define {{.+}}[[CALL00]]
 // CK31: define {{.+}}[[CALL01]]
 
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK32 --check-prefix CK32-32
+// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK32 --check-prefix CK32-32
+
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK32
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5383.region_id = weak constant i8 0
+// CK32: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5399.region_id = weak constant i8 0
+// CK32: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5416.region_id = weak constant i8 0
+// CK32: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 396]
+// CK32: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5432.region_id = weak constant i8 0
+// CK32: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: map_unused_var{{.*}}(
+void map_unused_var (){
+  float a;
+
+  // Region 00: default map type
+  // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK32-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK32-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK32-DAG: store float* [[VAR0]], float** [[CP0]]
+
+  // CK32: call void [[CALL00:@.+]](float* {{[^,]+}})
+  #pragma omp target map(a)
+  {}
+
+  // Region 01: non-default map type
+  // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK32-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
+  // CK32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
+  // CK32-DAG: store float* [[VAR1:%.+]], float** [[CBP1]]
+  // CK32-DAG: store float* [[VAR1]], float** [[CP1]]
+
+  // CK32: call void [[CALL01:@.+]](float* {{[^,]+}})
+  #pragma omp target map(to: a)
+  {}
+
+  // Region 02: non-scalar data type
+  // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK32-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [99 x float]**
+  // CK32-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [99 x float]**
+  // CK32-DAG: store [99 x float]* [[VAR2:%.+]], [99 x float]** [[CBP2]]
+  // CK32-DAG: store [99 x float]* [[VAR2]], [99 x float]** [[CP2]]
+
+  // CK32: call void [[CALL02:@.+]]([99 x float]* {{[^,]+}})
+  float arr[99];
+  #pragma omp target map(arr)
+  {}
+
+  // Region 03: used, but only in nested private region
+  // CK32-DAG: call i32 @__tgt_target_teams(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i32 0, i32 0)
+  // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK32-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK32-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to float**
+  // CK32-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to float**
+  // CK32-DAG: store float* [[VAR3:%.+]], float** [[CBP3]]
+  // CK32-DAG: store float* [[VAR3]], float** [[CP3]]
+
+  // CK32: call void [[CALL03:@.+]](float* {{[^,]+}})
+  #pragma omp target map(a)
+  #pragma omp teams private(a)
+  {
+    a++;
+  }
+}
+// CK32: define {{.+}}[[CALL00]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL01]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL02]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL03]]
+// CK32: call {{.*}} [[OUTLINE03:@\.omp_outlined\.[^ ]*]]
+// CK32: define {{.+}}[[OUTLINE03]]
+// CK32: alloca float
+
 #endif
 #endif