From 3994bafbc7b0bb1d4f21a2506be2a134aa6f886c Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Fri, 29 Jun 2018 16:16:00 +0000 Subject: [PATCH] [OPENMP, NVPTX] Sync threads before start ordered loops. Summary: Threads must be synchronized before starting ordered construct. Reviewers: grokos Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D48732 llvm-svn: 335987 --- openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu index 0e808df..1aab53e 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu @@ -240,12 +240,17 @@ public: // Process schedule. if (tnum == 1 || tripCount <= 1 || OrderedSchedule(schedule)) { + if (OrderedSchedule(schedule)) { + if (isSPMDMode()) + __syncthreads(); + else + __kmpc_barrier(loc, threadId); + } PRINT(LD_LOOP, "go sequential as tnum=%ld, trip count %lld, ordered sched=%d\n", (long)tnum, P64(tripCount), schedule); schedule = kmp_sched_static_chunk; chunk = tripCount; // one thread gets the whole loop - } else if (schedule == kmp_sched_runtime) { // process runtime omp_sched_t rtSched = currTaskDescr->GetRuntimeSched(); -- 2.7.4