[Openmp-commits] [openmp] r367024 - [OPENMP][NVPTX]Perform memory flush if number of threads to sync is 1 or less.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Thu Jul 25 08:02:28 2019
New Revision: 367024

URL: http://llvm.org/viewvc/llvm-project?rev=367024&view=rev
Log:
[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

Added:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c
Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu?rev=367024&r1=367023&r2=367024&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu Thu Jul 25 08:02:28 2019
@@ -62,6 +62,9 @@ EXTERN void __kmpc_barrier(kmp_Ident *lo
         // 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");
   }

Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c?rev=367024&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c Thu Jul 25 08:02:28 2019
@@ -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);
+}