[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 94983.
gtbercea added a comment.

Clean-up.


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,36 @@
       const llvm::function_ref<
           bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
                OpenMPClauseKind)> &Check) {
-    auto SI = Stack.rbegin();
-    auto SE = Stack.rend();
-
-    if (SI == SE)
+    if (Stack.empty())
       return false;
 
-    if (CurrentRegionOnly) {
-      SE = std::next(SI);
-    } else {
-      ++SI;
-    }
+    if (CurrentRegionOnly)
+      return checkMappableExprComponentListsForDeclAtLevel(VD, Stack.size() - 2, Check);
 
-    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;
-    }
+    for (unsigned I = Stack.size() - 1; I > 0; --I)
+      if (checkMappableExprComponentListsForDeclAtLevel(VD, I - 2, Check))
+        return true;
+    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) {
+    auto StartI = Stack.begin();
+    auto EndI = Stack.end();
+    if (std::distance(StartI, EndI) <= (int)(Level + 1))
+      return false;
+    std::advance(StartI, Level + 1);
+
+    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 +923,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


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