[OPENMP, NVPTX] Sync threads before start ordered loops.

[ABataev]

Summary

Threads must be synchronized before starting ordered construct.

[grokos]

The addition looks good. I've seen this piece of code in quite a few places, maybe we need to refactor it into a mini-function. But let's revisit that in the future.

libomptarget/deviceRTLs/nvptx/src/loop.cu

INLINE static void dispatch_init(kmp_Indent *loc, int32_t threadId,
      * them on, they merely affect which schedule we can legally choose for
      * various dynamic cases. (In paritcular, whether or not a stealing scheme
      * is legal).
      */
     schedule = SCHEDULE_WITHOUT_MODIFIERS(schedule);
 
     // 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();
       chunk = currTaskDescr->RuntimeChunkSize();
       switch (rtSched) {
       case omp_sched_static: {
         if (chunk > 0)
           schedule = kmp_sched_static_chunk;