[PATCH] D29905: [OpenMP] Pass argument to device kernel by reference when map is used.

[Gheorghe-Teodor Bercea via Phabricator via cfe-commits]

gtbercea updated this revision to Diff 94990.
gtbercea added a comment.

Fix for loop range.


Repository:
  rL LLVM

https://reviews.llvm.org/D29905

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


Index: test/OpenMP/target_map_codegen.cpp
===================================================================
--- test/OpenMP/target_map_codegen.cpp
+++ test/OpenMP/target_map_codegen.cpp
@@ -4756,3 +4756,20 @@
 }
 #endif
 #endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK30 -std=c++11 -fopenmp -S -emit-llvm -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s
+
+#ifdef CK30
+
+void target_maps_parallel_integer(int a){
+  int ParamToKernel = a;
+#pragma omp target map(tofrom: ParamToKernel)
+  {
+    ParamToKernel += 1;
+  }
+}
+
+// CK30: {{.*}}void @__omp_offloading_{{.*}}(i32* dereferenceable(4) %ParamToKernel) #0 {
+
+#endif
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -333,25 +333,37 @@
       const llvm::function_ref<
           bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
                OpenMPClauseKind)> &Check) {
-    auto SI = Stack.rbegin();
-    auto SE = Stack.rend();
+    if (Stack.empty())
+      return false;
+
+    if (CurrentRegionOnly)
+      return checkMappableExprComponentListsForDeclAtLevel(VD, Stack.size() - 1, Check);
+
+    if (Stack.size() > 2)
+      for (unsigned I = Stack.size() - 1; I > 0; --I)
+        if (checkMappableExprComponentListsForDeclAtLevel(VD, I - 1, Check))
+          return true;
+    return false;
+  }
 
-    if (SI == SE)
+  /// 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 (Stack.size() <= Level)
       return false;
 
-    if (CurrentRegionOnly) {
-      SE = std::next(SI);
-    } else {
-      ++SI;
-    }
+    auto StartI = Stack.begin();
+    std::advance(StartI, Level);
 
-    for (; SI != SE; ++SI) {
-      auto MI = SI->MappedExprComponents.find(VD);
-      if (MI != SI->MappedExprComponents.end())
-        for (auto &L : MI->second.Components)
-          if (Check(L, MI->second.Kind))
-            return true;
-    }
+    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;
   }
 
@@ -912,9 +924,8 @@
     bool IsVariableUsedInMapClause = false;
     bool IsVariableAssociatedWithSection = false;
 
-    DSAStack->checkMappableExprComponentListsForDecl(
-        D, /*CurrentRegionOnly=*/true,
-        [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
+    DSAStack->checkMappableExprComponentListsForDeclAtLevel(
+        D, Level + 1, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
                 MapExprComponents,
             OpenMPClauseKind WhereFoundClauseKind) {
           // Only the map clause information influences how a variable is


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D29905.94990.patch
Type: text/x-patch
Size: 3207 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170412/6c0c6845/attachment.bin>