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

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

gtbercea created this revision.

When a scalar argument is explicitly mapped, the device kernel needs to be passed that argument by reference.

This is a partial fix to this bug <https://bugs.llvm.org//show_bug.cgi?id=31920>.

The full fix requires data sharing support which will land in future patches.


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,24 @@
 }
 #endif
 #endif
+
+///==========================================================================///
+// RUN: %clang -DCK30 -fopenmp -S -emit-llvm -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s
+#ifdef CK30
+#include <stdio.h>
+
+void target_maps_parallel_integer(int a){
+  int ParamToKernel = a;
+  #pragma omp target map(tofrom:ParamToKernel)
+  {
+    printf("Print ParamToKernel = %d\n", ParamToKernel);
+    #pragma omp parallel
+    {
+      printf("Do nothing\n");
+    }
+  }
+}
+
+// CK30: define void @__omp_offloading_802_{{.*}}target_maps_parallel_integer{{.*}}l{{.*}}(i32* dereferenceable(4) %ParamToKernel) #0 {
+
+#endif
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -355,6 +355,27 @@
     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 = std::next(Stack.begin());
+    auto EndI = Stack.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(
@@ -912,9 +933,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.88243.patch
Type: text/x-patch
Size: 2683 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170213/2337dc42/attachment.bin>