[Openmp-commits] [PATCH] D62397: [OPENMP][NVPTX]Relax flush directive.

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added reviewers: grokos, gtbercea, kkwli0.
Herald added subscribers: jdoerfert, jfb, guansong.
Herald added a project: OpenMP.

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.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D62397

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


Index: libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
===================================================================
--- /dev/null
+++ libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
@@ -0,0 +1,33 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  int data, flag = 0;
+#pragma omp target parallel num_threads(64) map(to : data, flag)
+  {
+    if (omp_get_thread_num() == 0) {
+      /* Write to the data buffer that will be read by thread */
+      data = 42;
+/* Flush data to thread 1 and strictly order the write to data
+   relative to the write to the flag */
+#pragma omp flush(flag, data)
+      /* Set flag to release thread 1 */
+      flag = 1;
+/* Flush flag to ensure that thread 1 sees S-21 the change */
+#pragma omp flush(flag)
+    } else if (omp_get_thread_num() == 32) {
+/* Loop until we see the update to the flag */
+#pragma omp flush(flag, data)
+      while (flag < 1) {
+#pragma omp flush(flag, data)
+      }
+#pragma omp flush(flag, data)
+      // CHECK: data=42.
+      /* Values data will be 42, value of flag still undefined */
+      printf("data=%d.\n", data);
+    }
+  }
+  return 0;
+}
Index: libomptarget/deviceRTLs/nvptx/src/sync.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -130,7 +130,7 @@
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence_system();
+  __threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D62397.201257.patch
Type: text/x-patch
Size: 1649 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190524/f2ace6b6/attachment.bin>