[Openmp-commits] [PATCH] D86781: [LIBOMPTARGET]Do not try to optimize bases for the next parameters.

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added reviewers: jdoerfert, ye-luo, grokos.
Herald added a project: OpenMP.
ABataev requested review of this revision.
Herald added a subscriber: sstefan1.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D86781

Files:
  openmp/libomptarget/src/omptarget.cpp
  openmp/libomptarget/test/mapping/lambda_mapping.cpp


Index: openmp/libomptarget/test/mapping/lambda_mapping.cpp
===================================================================
--- /dev/null
+++ 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;
+}
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -1062,7 +1062,10 @@
       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());


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D86781.288590.patch
Type: text/x-patch
Size: 2429 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200828/5245a439/attachment-0001.bin>