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

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added a reviewer: grokos.
Herald added a subscriber: guansong.

Threads must be synchronized before starting ordered construct.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D48732

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


Index: libomptarget/deviceRTLs/nvptx/src/loop.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/loop.cu
+++ 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.153364.patch
Type: text/x-patch
Size: 883 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20180628/88ff0354/attachment.bin>