[LIBOMPTARGET]Do not try to optimize bases for the next parameters.

[ABataev]

Summary

PrivateArgumentManager shall immediately allocate firstprivates if they
are bases for the next parameters and the next paramaters rely on the
fact that the base musst be allocated already.

[grokos]

Looks good!

openmp/libomptarget/src/omptarget.cpp

1066

depend on

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1068

I have a quick question: Does front end order all arguments? For example, let's say a user write the following code:

#pragma omp target first private(a.p1[0:N], b, a.p2[0:N], c, a.p3[0:N])

where p1, p2 and p3 are members of a struct a. Then when passing to the runtime library, Args[I], Args[I+1] and Args[I+2] should be corresponding to a.p1[0:N], a.p2[0:N], a.p3[0:N], right? They'll not be interleaved by b and c.

openmp/libomptarget/test/mapping/lambda_mapping.cpp

11

Really interesting fix. I didn't realize that lambda function is captured as a first-private argument.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

1068

The order is important for correct passing of captured variables to the target regions (order of the arguments for the outlined target functions). For target regions it is not always defined by the order in the mappping/data-sharing clauses, but the order of captured variables, which can be different.
Yes, frontend already reorders maps to gather elements with the same bases. It is required for correct memory allocation of the partially copied structures.

openmp/libomptarget/test/mapping/lambda_mapping.cpp

11

In future we may try to implement the deepcopy of the record with data members having reference data types in the similar way.

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1068

I see. Thanks for the explanation.

openmp/libomptarget/src/omptarget.cpp

for (int32_t I = 0; I < ArgNum; ++I) {
     bool IsLast, IsHostPtr; // unused.
     if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       DP("Forwarding first-private value " DPxMOD " to the target construct\n",
          DPxPTR(HstPtrBase));
       TgtPtrBegin = HstPtrBase;
       TgtBaseOffset = 0;
     } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
-      const bool IsFirstPrivate = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
+      // Can be marked for optimization if the next argument(s) do(es) not
+      // depend from this one.
+      const bool IsFirstPrivate =
+          (I >= ArgNum - 1 || !(ArgTypes[I + 1] & OMP_TGT_MAPTYPE_MEMBER_OF));
       Ret = PrivateArgumentManager.addArg(HstPtrBegin, ArgSizes[I],
                                           TgtBaseOffset, IsFirstPrivate,
                                           TgtPtrBegin, TgtArgs.size());
       if (Ret != OFFLOAD_SUCCESS) {
         REPORT("Failed to process %sprivate argument " DPxMOD "\n",
                (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
         return OFFLOAD_FAIL;
       }

openmp/libomptarget/test/mapping/lambda_mapping.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <iostream>
+
+template <typename LOOP_BODY>
+inline void forall(int Begin, int End, LOOP_BODY LoopBody) {
+#pragma omp target parallel for schedule(static)
+  for (int I = Begin; I < End; ++I) {
+    LoopBody(I);
+  }
+}
+
+#define N (1000)
+
+//
+// Demonstration of the RAJA abstraction using lambdas
+// Requires data mapping onto the target section
+//
+int main() {
+  double A[N], B[N], C[N];
+
+  for (int I = 0; I < N; I++) {
+    A[I] = I + 1;
+    B[I] = -I;
+    C[I] = -9;
+  }
+
+#pragma omp target data map(tofrom : C [0:N]) map(to : A [0:N], B [0:N])
+  {
+    forall(0, N, [&](int I) { C[I] += A[I] + B[I]; });
+  }
+
+  int Fail = 0;
+  for (int I = 0; I < N; I++) {
+    if (C[I] != -8) {
+      std::cout << "Failed at " << I << " with val " << C[I] << std::endl;
+      Fail = 1;
+    }
+  }
+
+  // CHECK: Succeeded
+  if (Fail) {
+    std::cout << "Failed" << std::endl;
+  } else {
+    std::cout << "Succeeded" << std::endl;
+  }
+
+  return 0;
+}