[OPENMP][NVPTX]Relax flush directive.

[ABataev]

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.

[grokos]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

19

S-21? Typo?

[grokos]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

29

"Values" --> "Value of data". Also, why is the value of "flag" undefined at this point?

[ABataev]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

19

Yes, this is slightly reworked code from OpenMP Examples document and I forgot to remove this when copied the example.

29

as I understand, it happens because the compiler is free to reorder read/write ops in any order without explicitly provided ordering. That's why flag can be undefined here.

[jdoerfert]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

There is a race in this program.

After the flush in thread 0 happened, and after the flush prior to the while happened, both thread 0 and 32 will access flag without proper synchronization.

30

Couldn't we copy data to a new variable and check it on the host. Relying on printf here seems brittle.

[grokos]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

I thinks that's the purpose of this test: to make sure that, despite the race on flag, memory fencing is enforced for data.

[ABataev]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

Exactly, just like I said before this is slightly reworked example for flush directive from OpenMP examples.

30

Ok

[ABataev]

Address comments.

[jdoerfert]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

I thinks that's the purpose of this test: to make sure that, despite the race on flag, memory fencing is enforced for data.

Having a race on flag, especially since it is used in a control condition for the flush, makes the program as a whole UB. We cannot reliably test based on UB.

29

compiler is free to reorder read/write ops in any order without explicitly provided ordering.

Forgetting about the UB we have in this program for a moment, it is not true that a compiler is free to reorder arbitrary read/writes. I'm unsure what accesses you refer to here and how do they need to be reordered to lead to undefined values? (I could imagine this to be a side effect of the UB above which would make my point that this is not a reliable method to test things.)

[ABataev]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

This is slightly modified test from the official OpenMP document (see https://www.openmp.org/wp-content/uploads/openmp-examples-4.5.0.pdf, 8.1 The OpenMP Memory Model). If you think the test is not correct, you need to discuss it with OpenMP committee and create a ticket for this problem to fix the example.
The value of flag is undefined because of the data race and I think this was intended.

29

This assumption was wrong, actually, the value is undefined because of the race condition.

[ABataev]

Reworked test.

[jdoerfert]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

I will talk to the relevant people in the OpenMP ARB but that should not concern us here.
Whatever someone else thought would make a good test case should not impact us here.

UB is nothing to rely on and I fail to see why the race that causes it is an important part of the test case we need to keep.

[ABataev]

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

24

I agree. That's why I reworked the test and usу atomics for flag.

[jdoerfert]

Potential improvement inlined but this LGTM. Thanks for fixing the race problem!

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

34

Should we return !(out == 42)?

libomptarget/deviceRTLs/nvptx/src/sync.cu

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Flush
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence_system();
+  __threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Vote
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN int32_t __kmpc_warp_active_thread_mask() {
   PRINT0(LD_IO, "call __kmpc_warp_active_thread_mask\n");
   return __ACTIVEMASK();
 }

libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

+// 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 0;
+}