[Openmp-commits] [PATCH] D48732: [OPENMP, NVPTX] Sync threads before start ordered loops.

[Alexey Bataev via Phabricator via Openmp-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL335987: [OPENMP, NVPTX] Sync threads before start ordered loops. (authored by ABataev, committed by ).
Herald added a subscriber: llvm-commits.

Repository:
  rL LLVM

https://reviews.llvm.org/D48732

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


Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
@@ -240,12 +240,17 @@
 
     // 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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D48732.153506.patch
Type: text/x-patch
Size: 922 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20180629/1b96151c/attachment.bin>