diff options
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu | 2 | ||||
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c | 35 |
2 files changed, 36 insertions, 1 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu index d81aa8f0f3a..688420e0d7b 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -130,7 +130,7 @@ EXTERN void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid) { EXTERN void __kmpc_flush(kmp_Ident *loc) { PRINT0(LD_IO, "call kmpc_flush\n"); - __threadfence_system(); + __threadfence(); } //////////////////////////////////////////////////////////////////////////////// diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c new file mode 100644 index 00000000000..412538b6dd1 --- /dev/null +++ b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c @@ -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); +} |
