[OPENMP][NVPTX]Perform memory flush if number of threads to sync is 1 or less.

Summary:
According to the OpenMP standard, barrier operation must perform
implicit flush operation. Currently, if there is only one thread in the
team, barrier does not flush the memory. Patch fixes this problem.

Reviewers: grokos, gtbercea, kkwli0

Subscribers: guansong, jdoerfert, openmp-commits, caomhin

Tags: #openmp

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

git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@367024 91177308-0d34-0410-b5e6-96231b3b80d8

2 files changed, 40 insertions, 0 deletions

diff --git a/libomptarget/deviceRTLs/nvptx/src/sync.cu b/libomptarget/deviceRTLs/nvptx/src/sync.cu
index 688420e..4607d6a 100644
--- a/libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ b/libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -62,6 +62,9 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) {
         // Barrier #1 is for synchronization among active threads.
         named_sync(L1_BARRIER, threads);
       }
+    } else {
+      // Still need to flush the memory per the standard.
+      __kmpc_flush(loc_ref);
     } // numberOfActiveOMPThreads > 1
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
diff --git a/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c b/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c
new file mode 100644
index 0000000..7c70771
--- /dev/null
+++ b/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c
@@ -0,0 +1,37 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  int data, out, flag = 0;
+#pragma omp target teams num_teams(2) map(tofrom                               \
+                                          : out) map(to                        \
+                                                     : data, flag)             \
+    thread_limit(1)
+#pragma omp parallel num_threads(1)
+  {
+    if (omp_get_team_num() == 0) {
+      /* Write to the data buffer that will be read by thread in team 1 */
+      data = 42;
+/* Flush data to thread in team 1 */
+#pragma omp barrier
+      /* Set flag to release thread in team 1 */
+#pragma omp atomic write
+      flag = 1;
+    } else if (omp_get_team_num() == 1) {
+      /* Loop until we see the update to the flag */
+      int val;
+      do {
+#pragma omp atomic read
+        val = flag;
+      } while (val < 1);
+      out = data;
+#pragma omp barrier
+    }
+  }
+  // CHECK: out=42.
+  /* Value of out will be 42 */
+  printf("out=%d.\n", out);
+  return !(out == 42);
+}