[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

@@ -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();