[Openmp-commits] [openmp] 6aa7228 - [LIBOMPTARGET]Do not try to optimize bases for the next parameters.

[Alexey Bataev via Openmp-commits]

Author: Alexey Bataev
Date: 2020-08-28T15:46:31-04:00
New Revision: 6aa7228a629d81af78d4f701b7defb701f4b9283

URL: https://github.com/llvm/llvm-project/commit/6aa7228a629d81af78d4f701b7defb701f4b9283
DIFF: https://github.com/llvm/llvm-project/commit/6aa7228a629d81af78d4f701b7defb701f4b9283.diff

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

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.

Differential Revision: https://reviews.llvm.org/D86781

Added: 
    openmp/libomptarget/test/mapping/lambda_mapping.cpp

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index e20a2c9d4b24..c4f781d46959 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -1062,7 +1062,10 @@ int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
       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 on 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());

diff  --git a/openmp/libomptarget/test/mapping/lambda_mapping.cpp b/openmp/libomptarget/test/mapping/lambda_mapping.cpp
new file mode 100644
index 000000000000..23671ebb7d68
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/lambda_mapping.cpp
@@ -0,0 +1,53 @@
+// 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;
+}