diff options
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu | 3 | ||||
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c | 37 |
2 files changed, 40 insertions, 0 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu index 688420e0d7b..4607d6a7d8a 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -62,6 +62,9 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) { // Barrier #1 is for synchronization among active threads. named_sync(L1_BARRIER, threads); } + } else { + // Still need to flush the memory per the standard. + __kmpc_flush(loc_ref); } // numberOfActiveOMPThreads > 1 PRINT0(LD_SYNC, "completed kmpc_barrier\n"); } diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c new file mode 100644 index 00000000000..7c707718e13 --- /dev/null +++ b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c @@ -0,0 +1,37 @@ +// 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 teams num_teams(2) map(tofrom \ + : out) map(to \ + : data, flag) \ + thread_limit(1) +#pragma omp parallel num_threads(1) + { + if (omp_get_team_num() == 0) { + /* Write to the data buffer that will be read by thread in team 1 */ + data = 42; +/* Flush data to thread in team 1 */ +#pragma omp barrier + /* Set flag to release thread in team 1 */ +#pragma omp atomic write + flag = 1; + } else if (omp_get_team_num() == 1) { + /* 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 barrier + } + } + // CHECK: out=42. + /* Value of out will be 42 */ + printf("out=%d.\n", out); + return !(out == 42); +} |
