[OpenMP] Fix target map for unused variables

[jdenny]

Summary

Without this patch, each of the following map clauses doesn't map
its variable into the target region because the variable is unused in
the target region, as discussed in D65835#1624669:

#pragma omp target map(a)
{}

#pragma omp target map(a)
#pragma omp teams private(a)
{
  a++;
}

This patch fixes that by marking all map clause variables for
capturing. That means the capturing analysis now sometimes runs on a
capture region within a combined construct, so this patch adjusts the
analysis to be precise about how many capture regions remain in a
combined construct. Otherwise, existing tests break.

[ABataev]

Do we really need to map such variables? According to standard, "The map clause specifies how an original list item is mapped from the current task’s data environment to a corresponding list item in the device data environment of the device identified by the construct.", i.e. it does not map the variable, but "specifies" how to map it.

[jdenny]

In D66247#1630196, @ABataev wrote:

Do we really need to map such variables? According to standard, "The map clause specifies how an original list item is mapped from the current task’s data environment to a corresponding list item in the device data environment of the device identified by the construct.", i.e. it does not map the variable, but "specifies" how to map it.

What you're asking for now contradicts what you asked for at https://reviews.llvm.org/D65835#1624659 when I asked about the same example of target map(a) enclosing teams private(a).

[ABataev]

In D66247#1630245, @jdenny wrote:

In D66247#1630196, @ABataev wrote:

Do we really need to map such variables? According to standard, "The map clause specifies how an original list item is mapped from the current task’s data environment to a corresponding list item in the device data environment of the device identified by the construct.", i.e. it does not map the variable, but "specifies" how to map it.

What you're asking for now contradicts what you asked for at https://reviews.llvm.org/D65835#1624659 when I asked about the same example of target map(a) enclosing teams private(a).

Seems to me I did not expressed my idea absolutely correctly. I meant that if the variable is really mapped, then the maptype must be generated in all cases. If we do not need to map the variable, just like in these case, there definitely should not be maptype.

[jdenny]

In D66247#1630262, @ABataev wrote:

In D66247#1630245, @jdenny wrote:

In D66247#1630196, @ABataev wrote:

Do we really need to map such variables? According to standard, "The map clause specifies how an original list item is mapped from the current task’s data environment to a corresponding list item in the device data environment of the device identified by the construct.", i.e. it does not map the variable, but "specifies" how to map it.

What you're asking for now contradicts what you asked for at https://reviews.llvm.org/D65835#1624659 when I asked about the same example of target map(a) enclosing teams private(a).

Seems to me I did not expressed my idea absolutely correctly. I meant that if the variable is really mapped, then the maptype must be generated in all cases. If we do not need to map the variable, just like in these case, there definitely should not be maptype.

Then I still need an example where D65835 doesn't behave correctly.

[ABataev]

In D66247#1630287, @jdenny wrote:

In D66247#1630262, @ABataev wrote:

In D66247#1630245, @jdenny wrote:

In D66247#1630196, @ABataev wrote:

Do we really need to map such variables? According to standard, "The map clause specifies how an original list item is mapped from the current task’s data environment to a corresponding list item in the device data environment of the device identified by the construct.", i.e. it does not map the variable, but "specifies" how to map it.

What you're asking for now contradicts what you asked for at https://reviews.llvm.org/D65835#1624659 when I asked about the same example of target map(a) enclosing teams private(a).

Seems to me I did not expressed my idea absolutely correctly. I meant that if the variable is really mapped, then the maptype must be generated in all cases. If we do not need to map the variable, just like in these case, there definitely should not be maptype.

Then I still need an example where D65835 doesn't behave correctly.

Try map(a) firstprivate(a) defaultmap(scalar:tofrom), where a is int, for example. The variable must be mapped as tofrom in this case but, most probably, will be mapped as to.

[jdenny]

In D66247#1630321, @ABataev wrote:

Try map(a) firstprivate(a) defaultmap(scalar:tofrom), where a is int, for example. The variable must be mapped as tofrom in this case but, most probably, will be mapped as to.

Map type is 35=0x23, which has tofrom.

[ABataev]

In D66247#1630333, @jdenny wrote:

In D66247#1630321, @ABataev wrote:

Try map(a) firstprivate(a) defaultmap(scalar:tofrom), where a is int, for example. The variable must be mapped as tofrom in this case but, most probably, will be mapped as to.

Map type is 35=0x23, which has tofrom.

Yes, just realized that, defaultmap does not affect explicit firstprivates. Then just check map(a) firstprivate(a) for int128 type. Check that it still has tofrom mapping. If so, then probably we already can handle such combinations correctly in the codegen. Also, test it for other mapping types, like alloc, from, to.

[jdenny]

In D66247#1630352, @ABataev wrote:

Yes, just realized that, defaultmap does not affect explicit firstprivates. Then just check map(a) firstprivate(a) for int128 type. Check that it still has tofrom mapping. If so, then probably we already can handle such combinations correctly in the codegen. Also, test it for other mapping types, like alloc, from, to.

All work as expected.

[ABataev]

In D66247#1630398, @jdenny wrote:

In D66247#1630352, @ABataev wrote:

Yes, just realized that, defaultmap does not affect explicit firstprivates. Then just check map(a) firstprivate(a) for int128 type. Check that it still has tofrom mapping. If so, then probably we already can handle such combinations correctly in the codegen. Also, test it for other mapping types, like alloc, from, to.

All work as expected.

Then just add the tests for as many different combinations as possible to be sure that we did not miss anything.

clang/include/clang/Sema/ScopeInfo.h

public:
 
   /// The implicit parameter for the captured variables.
   ImplicitParamDecl *ContextParam;
 
   /// The kind of captured region.
   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;
   }
 
   ~CapturedRegionScopeInfo() override;
 
   /// A descriptive name for the kind of captured region this is.
   StringRef getRegionName() const {
     switch (CapRegionKind) {

clang/include/clang/Sema/Sema.h

private:
   void InitDataSharingAttributesStack();
   void DestroyDataSharingAttributesStack();
   ExprResult
   VerifyPositiveIntegerConstantInClause(Expr *Op, OpenMPClauseKind CKind,
                                         bool StrictlyPositive = true);
   /// 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();
 
   /// Pop OpenMP function region for non-capturing function.
   void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
 
   /// Check whether we're allowed to call Callee from the current function.
   void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee);

clang/lib/Sema/Sema.cpp

 IdentifierInfo *Sema::getFloat128Identifier() const {
   if (!Ident___float128)
     Ident___float128 = &Context.Idents.get("__float128");
   return Ident___float128;
 }
 
 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);
 }
 
 CapturedRegionScopeInfo *Sema::getCurCapturedRegion() {
   if (FunctionScopes.empty())
     return nullptr;
 
   return dyn_cast<CapturedRegionScopeInfo>(FunctionScopes.back());

clang/lib/Sema/SemaExpr.cpp

         if (RSI->CapRegionKind == CR_OpenMP) {
           bool IsOpenMPPrivateDecl = isOpenMPPrivateDecl(Var, RSI->OpenMPLevel);
           auto IsTargetCap = !IsOpenMPPrivateDecl &&
                              isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel);
           // When we detect target captures we are looking from inside the
           // 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;
             DeclRefType = DeclRefType.getUnqualifiedType();
             CaptureType = Context.getLValueReferenceType(DeclRefType);
             break;
           }
         }

clang/lib/Sema/SemaOpenMP.cpp

DVarPrivate = DSAStack->hasDSA(D, isOpenMPPrivate,
                                    DSAStack->isClauseParsingMode());
     if (DVarPrivate.CKind != OMPC_unknown ||
         (VD && DSAStack->getDefaultDSA() == DSA_none))
       return VD ? VD : cast<VarDecl>(DVarPrivate.PrivateCopy->getDecl());
   }
   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()))
     DSAStack->loopInit();
 }
 
 bool Sema::isOpenMPPrivateDecl(const ValueDecl *D, unsigned Level) const {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
StmtResult Sema::ActOnOpenMPRegionEnd(StmtResult S,
   }
 
   SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
   getOpenMPCaptureRegions(CaptureRegions, DSAStack->getCurrentDirective());
   OMPOrderedClause *OC = nullptr;
   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()) &&
         Clause->getClauseKind() == OMPC_in_reduction) {
       // Capture taskgroup task_reduction descriptors inside the tasking regions
       // with the corresponding in_reduction items.
       auto *IRC = cast<OMPInReductionClause>(Clause);
       for (Expr *E : IRC->taskgroup_descriptors())
if (isOpenMPPrivate(Clause->getClauseKind()) ||
       }
     }
     if (Clause->getClauseKind() == OMPC_schedule)
       SC = cast<OMPScheduleClause>(Clause);
     else if (Clause->getClauseKind() == OMPC_ordered)
       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
   // specified.
   if (SC &&
       (SC->getFirstScheduleModifier() == OMPC_SCHEDULE_MODIFIER_nonmonotonic ||
        SC->getSecondScheduleModifier() ==
            OMPC_SCHEDULE_MODIFIER_nonmonotonic) &&
if (ThisCaptureRegion != OMPD_unknown) {
             CaptureRegion == OMPD_unknown) {
           if (auto *DS = cast_or_null<DeclStmt>(C->getPreInitStmt())) {
             for (Decl *D : DS->decls())
               MarkVariableReferenced(D->getLocation(), cast<VarDecl>(D));
           }
         }
       }
     }
+    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());
   }
   return SR;
 }
 
 static bool checkCancelRegion(Sema &SemaRef, OpenMPDirectiveKind CurrentRegion,

clang/test/OpenMP/target_map_codegen.cpp

void explicit_maps_single (int ii){
   {
    a++;
   }
 }
 // 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