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

[gtbercea]

Summary

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.

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

[gtbercea]

Update test.

[ABataev]

lib/Sema/SemaOpenMP.cpp

358–360

Could you join these 2 functions into one?

[gtbercea]

lib/Sema/SemaOpenMP.cpp

358–360

Since SI and StarI use different types of iterations the functions cannot be merged.

[ABataev]

lib/Sema/SemaOpenMP.cpp

358–360

I still believe you can reuse this new code:

bool checkMappableExprComponentListsForDecl(
    ValueDecl *VD, bool CurrentRegionOnly,
    const llvm::function_ref<
        bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
             OpenMPClauseKind)> &Check) {
  if (Stack.empty())
    return false;

  if (CurrentRegionOnly)
    return checkMappableExprComponentListsForDeclAtLevel(VD, Stack.size() - 1, Check);
  
  for (unsigned I = Stack.size(); I > 0; --I)
    if (checkMappableExprComponentListsForDeclAtLevel(VD, I - 1, Check))
      return true; 
  return false;
}

or something like this.

[gtbercea]

Refactor code.

[gtbercea]

Pass correct levels. Refactor code.

[ABataev]

lib/Sema/SemaOpenMP.cpp

355

Are you sure this is correct? If I is 1, I-2 will give you 0xFFFFFFFF. Please, check everything one more time

367

Please, remove this commented code

[gtbercea]

Clean-up.

[gtbercea]

Refactor code.

[gtbercea]

Fix for loop range.

[gtbercea]

Update check before loop.

[ABataev]

LG

[Hahnfeld]

In D34888#799576, @gtbercea wrote:

Does this also include the fixes in the following revision?

https://reviews.llvm.org/D29905

Sorry, I wasn't aware of this revision and thought that it had long been committed. I just verified that the bug referenced in the summary is also fixed by my patch in D34888. However, I can't comment on whether this patch is still needed. Sorry for the conflicts if yes...

You probably should commit your patches earlier, you currently have 10 accepted revisions that have not yet been committed. This will also avoid complicated rebases and so on.

[gtbercea]

In D29905#800416, @Hahnfeld wrote:

In D34888#799576, @gtbercea wrote:

Does this also include the fixes in the following revision?

https://reviews.llvm.org/D29905

Sorry, I wasn't aware of this revision and thought that it had long been committed. I just verified that the bug referenced in the summary is also fixed by my patch in D34888. However, I can't comment on whether this patch is still needed. Sorry for the conflicts if yes...

You probably should commit your patches earlier, you currently have 10 accepted revisions that have not yet been committed. This will also avoid complicated rebases and so on.

No problem at all! I am trying to commit the rest of the patches. Currently waiting on one more patch to get approved that will unlock the rest. I do encourage you to review/comment on it: https://reviews.llvm.org/D34784 :)

[gtbercea]

Already covered by D34888

lib/Sema/SemaOpenMP.cpp

bool checkMappableExprComponentListsForDecl(
 
     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;
     }
     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(
       ValueDecl *VD,
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
       OpenMPClauseKind WhereFoundClauseKind) {
     assert(Stack.size() > 1 &&
            "Not expecting to retrieve components from a empty stack!");
if (Ty->isReferenceType())
       Ty = Ty->castAs<ReferenceType>()->getPointeeType();
 
     // Locate map clauses and see if the variable being captured is referred to
     // in any of those clauses. Here we only care about variables, not fields,
     // because fields are part of aggregates.
     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
           // captured. E.g. is_device_ptr does not require changing the default
           // behavior.
           if (WhereFoundClauseKind != OMPC_map)
             return false;
 

test/OpenMP/target_map_codegen.cpp

 void explicit_maps_member_pointer_references(SSA *sap) {
   double *d;
   SSA sa(d);
   SSB sb(sap);
   sb.foo();
 }
 #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