[OpenMP] Fix mapping of scalars for combined directives

Combined directives like 'target parallel' have two captured statements.
Sema has to check the right one from the right direction.

Previously, Sema::IsOpenMPCapturedByRef would return false for mapped
scalars on combined directives. This results in a wrong signature of
the outlined function which triggers an assertion:
void llvm::CallInst::init(llvm::FunctionType *, llvm::Value *, ArrayRef<llvm::Value *>, ArrayRef<OperandBundleDef>, const llvm::Twine &): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.

Fixes PR30975 (and PR31985). New function was taken from clang-ykt.

Differential Revision: https://reviews.llvm.org/D34888

llvm-svn: 306956

clang/lib/Sema/SemaOpenMP.cpp

@@ -412,6 +412,30 @@ public:
     return false;
   }
 
+  /// Do the check specified in \a Check to all component lists at a given level
+  /// and return true if any issue is found.
+  bool checkMappableExprComponentListsForDeclAtLevel(
+      ValueDecl *VD, unsigned Level,
+      const llvm::function_ref<
+          bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
+               OpenMPClauseKind)> &Check) {
+    if (isStackEmpty())
+      return false;
+
+    auto StartI = Stack.back().first.begin();
+    auto EndI = Stack.back().first.end();
+    if (std::distance(StartI, EndI) <= (int)Level)
+      return false;
+    std::advance(StartI, Level);
+
+    auto MI = StartI->MappedExprComponents.find(VD);
+    if (MI != StartI->MappedExprComponents.end())
+      for (auto &L : MI->second.Components)
+        if (Check(L, MI->second.Kind))
+          return true;
+    return false;
+  }
+
   /// Create a new mappable expression component list associated with a given
   /// declaration and initialize it with the provided list of components.
   void addMappableExpressionComponents(
@@ -994,9 +1018,8 @@ bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, unsigned Level) {
     bool IsVariableUsedInMapClause = false;
     bool IsVariableAssociatedWithSection = false;
 
-    DSAStack->checkMappableExprComponentListsForDecl(
-        D, /*CurrentRegionOnly=*/true,
-        [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
+    DSAStack->checkMappableExprComponentListsForDeclAtLevel(
+        D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
                 MapExprComponents,
             OpenMPClauseKind WhereFoundClauseKind) {
           // Only the map clause information influences how a variable is

clang/test/OpenMP/target_map_codegen.cpp

@@ -1056,6 +1056,9 @@ void implicit_maps_template_type_capture (int a){
 // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
 // CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
+// CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK19: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
+
 // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
 // CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
@@ -1194,6 +1197,28 @@ void explicit_maps_single (int ii){
     ++a;
   }
 
+  // Map of a scalar in nested region.
+  int b = a;
+
+  // Region 00n
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00n]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00n]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]]
+
+  // CK19: call void [[CALL00n:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(alloc:b)
+  #pragma omp parallel
+  {
+    ++b;
+  }
+
   // Map of an array.
   int arra[100];
 
@@ -2388,6 +2413,7 @@ void explicit_maps_single (int ii){
 }
 
 // CK19: define {{.+}}[[CALL00]]
+// CK19: define {{.+}}[[CALL00n]]
 // CK19: define {{.+}}[[CALL01]]
 // CK19: define {{.+}}[[CALL02]]
 // CK19: define {{.+}}[[CALL03]]