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
llvm-svn: 364572
EXTERN void __kmpc_flush(kmp_Ident *loc) {
PRINT0(LD_IO, "call kmpc_flush\n");
- __threadfence_system();
+ __threadfence();
}
////////////////////////////////////////////////////////////////////////////////
--- /dev/null
+// 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);
+}