[Openmp-commits] [openmp] r364572 - [OPENMP][NVPTX]Relax flush directive.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Thu Jun 27 11:33:09 2019
New Revision: 364572

URL: http://llvm.org/viewvc/llvm-project?rev=364572&view=rev
Log:
[OPENMP][NVPTX]Relax flush directive.

Summary:
According to the OpenMP standard, flush  makes a thread’s temporary view of memory consistent with memory and enforces an order on the memory operations of the variables explicitly specified or implied.

According to the Cuda toolkit documentation (https://docs.nvidia.com/cuda/archive/8.0/cuda-c-programming-guide/index.html#memory-fence-functions), __threadfence() functions provides required functionality.

__threadfence_system() also provides required functionality, but it also
includes some extra functionality, like synchronization of page-locked
host memory, synchronization for the host, etc. It is not required per
the standard and we can use more relaxed version of memory fence
operation.

Reviewers: grokos, gtbercea, kkwli0

Subscribers: guansong, jfb, jdoerfert, openmp-commits, caomhin

Tags: #openmp

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

Added:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/flush.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=364572&r1=364571&r2=364572&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu Thu Jun 27 11:33:09 2019
@@ -130,7 +130,7 @@ EXTERN void __kmpc_end_single(kmp_Ident
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence_system();
+  __threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////

Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c?rev=364572&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c Thu Jun 27 11:33:09 2019
@@ -0,0 +1,35 @@
+// 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 parallel num_threads(64) map(tofrom                         \
+                                                : out, flag) map(to            \
+                                                                 : data)
+  {
+    if (omp_get_thread_num() == 0) {
+      /* Write to the data buffer that will be read by thread */
+      data = 42;
+/* Flush data to thread 32 */
+#pragma omp flush(data)
+      /* Set flag to release thread 32 */
+#pragma omp atomic write
+      flag = 1;
+    } else if (omp_get_thread_num() == 32) {
+      /* 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 flush(out)
+    }
+  }
+  // CHECK: out=42.
+  /* Value of out will be 42 */
+  printf("out=%d.\n", out);
+  return !(out == 42);
+}