[Openmp-commits] [openmp] r335987 - [OPENMP, NVPTX] Sync threads before start ordered loops.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Fri Jun 29 09:16:00 2018
New Revision: 335987

URL: http://llvm.org/viewvc/llvm-project?rev=335987&view=rev
Log:
[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

Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu?rev=335987&r1=335986&r2=335987&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu Fri Jun 29 09:16:00 2018
@@ -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();