[PATCH] D34888: [OpenMP] Fix mapping of scalars for combined directives

[Jonas Hahnfeld via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL306956: [OpenMP] Fix mapping of scalars for combined directives (authored by Hahnfeld).

Changed prior to commit:
  https://reviews.llvm.org/D34888?vs=104863&id=104986#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D34888

Files:
  cfe/trunk/lib/Sema/SemaOpenMP.cpp
  cfe/trunk/test/OpenMP/target_map_codegen.cpp


Index: cfe/trunk/lib/Sema/SemaOpenMP.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp
@@ -412,6 +412,30 @@
     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 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
Index: cfe/trunk/test/OpenMP/target_map_codegen.cpp
===================================================================
--- cfe/trunk/test/OpenMP/target_map_codegen.cpp
+++ cfe/trunk/test/OpenMP/target_map_codegen.cpp
@@ -1056,6 +1056,9 @@
 // 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 @@
     ++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 @@
 }
 
 // CK19: define {{.+}}[[CALL00]]
+// CK19: define {{.+}}[[CALL00n]]
 // CK19: define {{.+}}[[CALL01]]
 // CK19: define {{.+}}[[CALL02]]
 // CK19: define {{.+}}[[CALL03]]


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D34888.104986.patch
Type: text/x-patch
Size: 3851 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170701/ce37cee6/attachment.bin>