Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu =================================================================== --- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -62,6 +62,9 @@ // 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"); } Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c =================================================================== --- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c +++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c @@ -0,0 +1,37 @@ +// RUN: %compile-run-and-check + +#include +#include + +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); +}