[OPENMP, NVPTX] Fixes for NVPTX RTL

[ABataev]

Summary

Patch fixes several problems in the implementation of NVPTX RTL.

1. Detection of the last iteration for loops with static scheduling, no chunks.
2. Fixes reductions for the serialized parallel constructs.
3. Fixes handling of the barriers.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/sync.cu

36–38 ↗

(On Diff #152462)

Are you moving these functions for a reason? I think they are defined as extern, so there's no need to have them before calling them.

(Moving code makes reading patches very difficult, I think the only change in this file is in __kmpc_barrier, right?)

[ABataev]

libomptarget/deviceRTLs/nvptx/src/sync.cu

36–38 ↗

(On Diff #152462)

Right, only __kmpc_[cancel_]barrier has changed. But it uses these generic|spmd versions and the compiler cannot see them because they are defined/declared later.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/sync.cu

36–38 ↗

(On Diff #152462)

They are declared in interface.h which should be included from omptarget-nvptx.h, no?

[ABataev]

libomptarget/deviceRTLs/nvptx/src/sync.cu

36–38 ↗

(On Diff #152462)

Ah, yes. Ok, I'll restore the original order.

[ABataev]

Restored original order of the barrier functions

[Hahnfeld]

Sounds good to me, I'll wait for @grokos to take a look.

[grokos]

Looks good to me, thanks for submitting those fixes!

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

if (entityId < leftOver) {
       chunk++;
       lb = lb + entityId * chunk;
     } else {
       lb = lb + entityId * chunk + leftOver;
     }
 
     T inputUb = ub;
     ub = lb + chunk - 1; // Clang uses i <= ub
-    last = ub == inputUb;
+    last = lb <= inputUb && inputUb <= ub;
     stride = loopSize; // make sure we only do 1 chunk per warp
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for Static Init
 
   INLINE static void for_static_init(int32_t schedtype, int32_t *plastiter,
                                      T *plower, T *pupper, ST *pstride,

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/reduction.cu

 
 INLINE
 int32_t nvptx_parallel_reduce_nowait(int32_t global_tid, int32_t num_vars,
                                      size_t reduce_size, void *reduce_data,
                                      kmp_ShuffleReductFctPtr shflFct,
                                      kmp_InterWarpCopyFctPtr cpyFct,
                                      bool isSPMDExecutionMode,
                                      bool isRuntimeUninitialized = false) {
+  uint32_t BlockThreadId = GetLogicalThreadIdInBlock();
+  uint32_t NumThreads = GetNumberOfOmpThreads(
+      BlockThreadId, isSPMDExecutionMode, isRuntimeUninitialized);
+  if (NumThreads == 1)
+    return 1;
   /*
    * This reduce function handles reduction within a team. It handles
    * parallel regions in both L1 and L2 parallelism levels. It also
    * supports Generic, SPMD, and NoOMP modes.
    *
    * 1. Reduce within a warp.
    * 2. Warp master copies value to warp 0 via shared memory.
    * 3. Warp 0 reduces to a single value.
    * 4. The reduced value is available in the thread that returns 1.
    */
 
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
-  uint32_t BlockThreadId = GetLogicalThreadIdInBlock();
-  uint32_t NumThreads = GetNumberOfOmpThreads(
-      BlockThreadId, isSPMDExecutionMode, isRuntimeUninitialized);
   uint32_t WarpsNeeded = (NumThreads + WARPSIZE - 1) / WARPSIZE;
   uint32_t WarpId = BlockThreadId / WARPSIZE;
 
   // Volta execution model:
   // For the Generic execution mode a parallel region either has 1 thread and
   // beyond that, always a multiple of 32. For the SPMD execution mode we may
   // have any number of threads.
   if ((NumThreads % WARPSIZE == 0) || (WarpId < WarpsNeeded - 1))
else if (!(Liveness & (Liveness + 1))) // Partial warp but contiguous lanes
     gpu_irregular_warp_reduce(reduce_data, shflFct,
                               /*LaneCount=*/__popc(Liveness),
                               /*LaneId=*/GetThreadIdInBlock() % WARPSIZE);
   else if (!isRuntimeUninitialized) // Dispersed lanes. Only threads in L2
                                     // parallel region may enter here; return
                                     // early.
     return gpu_irregular_simd_reduce(reduce_data, shflFct);
 
-  uint32_t BlockThreadId = GetLogicalThreadIdInBlock();
-  uint32_t NumThreads = GetNumberOfOmpThreads(
-      BlockThreadId, isSPMDExecutionMode, isRuntimeUninitialized);
-
   // When we have more than [warpsize] number of threads
   // a block reduction is performed here.
   //
   // Only L1 parallel region can enter this if condition.
   if (NumThreads > WARPSIZE) {
     uint32_t WarpsNeeded = (NumThreads + WARPSIZE - 1) / WARPSIZE;
     // Gather all the reduced values from each warp
     // to the first warp.

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu

 ////////////////////////////////////////////////////////////////////////////////
 
 // a team is a block: we can use CUDA native synchronization mechanism
 // FIXME: what if not all threads (warps) participate to the barrier?
 // We may need to implement it differently
 
 EXTERN int32_t __kmpc_cancel_barrier(kmp_Indent *loc_ref, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_cancel_barrier\n");
-  __syncthreads();
+  __kmpc_barrier(loc_ref, tid);
   PRINT0(LD_SYNC, "completed kmpc_cancel_barrier\n");
   return 0;
 }
 
 EXTERN void __kmpc_barrier(kmp_Indent *loc_ref, int32_t tid) {
+  if (isSPMDMode()) {
+    __kmpc_barrier_simple_spmd(loc_ref, tid);
+  } else if (isRuntimeUninitialized()) {
+    __kmpc_barrier_simple_generic(loc_ref, tid);
+  } else {
     tid = GetLogicalThreadIdInBlock();
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(tid);
     if (!currTaskDescr->InL2OrHigherParallelRegion()) {
       int numberOfActiveOMPThreads =
           GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized());
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
       // On Volta and newer architectures we require that all lanes in
       // a warp (at least, all present for the kernel launch) participate in the
       // barrier.  This is enforced when launching the parallel region.  An
       // exception is when there are < WARPSIZE workers.  In this case only 1
       // worker is started, so we don't need a barrier.
       if (numberOfActiveOMPThreads > 1) {
 #endif
         // The #threads parameter must be rounded up to the WARPSIZE.
         int threads =
             WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);
 
         PRINT(LD_SYNC,
               "call kmpc_barrier with %d omp threads, sync parameter %d\n",
               numberOfActiveOMPThreads, threads);
         // Barrier #1 is for synchronization among active threads.
         named_sync(L1_BARRIER, threads);
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
       } // numberOfActiveOMPThreads > 1
 #endif
     }
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
+}
 
 // Emit a simple barrier call in SPMD mode.  Assumes the caller is in an L0
 // parallel region and that all worker threads participate.
 EXTERN void __kmpc_barrier_simple_spmd(kmp_Indent *loc_ref, int32_t tid) {
   PRINT0(LD_SYNC, "call kmpc_barrier_simple_spmd\n");
   __syncthreads();
   PRINT0(LD_SYNC, "completed kmpc_barrier_simple_spmd\n");
 }