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

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev updated this revision to Diff 206900.
ABataev added a comment.
Herald added a subscriber: jfb.

Updated test.


Repository:
  rOMP OpenMP

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D62398/new/

https://reviews.llvm.org/D62398

Files:
  libomptarget/deviceRTLs/nvptx/src/sync.cu
  libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c


Index: libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c
===================================================================
--- /dev/null
+++ 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);
+}
Index: libomptarget/deviceRTLs/nvptx/src/sync.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -62,6 +62,9 @@
         // 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");
   }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D62398.206900.patch
Type: text/x-patch
Size: 1849 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190627/0f54822d/attachment.bin>