[OPENMP][NVPTX]Improve number of threads counter, NFC.

[ABataev]

Summary

Patch improves performance of the full runtime mode by moving
number-of-threads counter to the shared memory. It also allows to save
global memory.

[grokos]

LGTM

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

EXTERN void omp_set_num_threads(int num) {
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);
     currTaskDescr->NThreads() = num;
   }
 }
 
 EXTERN int omp_get_num_threads(void) {
-  bool isSPMDExecutionMode = isSPMDMode();
-  int tid = GetLogicalThreadIdInBlock(isSPMDExecutionMode);
-  int rc = GetNumberOfOmpThreads(tid, isSPMDExecutionMode);
+  int rc = GetNumberOfOmpThreads(isSPMDMode());
   PRINT(LD_IO, "call omp_get_num_threads() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_max_threads(void) {
   if (parallelLevel[GetWarpId()] > 0)
     // We're already in parallel region.
     return 1; // default is 1 thread avail

 }
 
 EXTERN int omp_get_ancestor_thread_num(int level) {
   if (isSPMDMode())
     return level == 1 ? GetThreadIdInBlock() : 0;
   int rc = -1;
   // If level is 0 or all parallel regions are not active - return 0.
   unsigned parLevel = parallelLevel[GetWarpId()];
-  if (level == 0 || (level > 0 && parLevel < OMP_ACTIVE_PARALLEL_LEVEL &&
-                     level <= parLevel)) {
-    rc = 0;
-  } else if (level > 0) {
+  if (level == 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL) {
     int totLevel = omp_get_level();
     if (level <= totLevel) {
       omptarget_nvptx_TaskDescr *currTaskDescr =
           getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);
       int steps = totLevel - level;
       PRINT(LD_IO, "backtrack %d steps\n", steps);
       ASSERT0(LT_FUSSY, currTaskDescr,
               "do not expect fct to be called in a non-active thread");
       do {
         if (DON(LD_IOD)) {
           // print current state
           omp_sched_t sched = currTaskDescr->GetRuntimeSched();
           PRINT(LD_ALL,
                 "task descr %s %d: %s, in par %d, rt sched %d,"
                 " chunk %" PRIu64 "; tid %d, tnum %d, nthreads %d\n",
                 "ancestor", steps,
                 (currTaskDescr->IsParallelConstruct() ? "par" : "task"),
                 (int)currTaskDescr->InParallelRegion(), (int)sched,
                 currTaskDescr->RuntimeChunkSize(),
-                (int)currTaskDescr->ThreadId(),
-                (int)currTaskDescr->ThreadsInTeam(),
+                (int)currTaskDescr->ThreadId(), (int)threadsInTeam,
                 (int)currTaskDescr->NThreads());
         }
 
         if (currTaskDescr->IsParallelConstruct()) {
           // found the level
           if (!steps) {
             rc = currTaskDescr->ThreadId();
             break;
           }
           steps--;
         }
         currTaskDescr = currTaskDescr->GetPrevTaskDescr();
       } while (currTaskDescr);
       ASSERT0(LT_FUSSY, !steps, "expected to find all steps");
     }
+  } else if (level == 0 ||
+             (level > 0 && parLevel < OMP_ACTIVE_PARALLEL_LEVEL &&
+              level <= parLevel) ||
+             (level > 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL &&
+              level <= (parLevel - OMP_ACTIVE_PARALLEL_LEVEL))) {
+    rc = 0;
   }
   PRINT(LD_IO, "call omp_get_ancestor_thread_num(level %d) returns %d\n", level,
         rc)
   return rc;
 }
 
 EXTERN int omp_get_team_size(int level) {
   if (isSPMDMode())
     return level == 1 ? GetNumberOfThreadsInBlock() : 1;
   int rc = -1;
   unsigned parLevel = parallelLevel[GetWarpId()];
   // If level is 0 or all parallel regions are not active - return 1.
-  if (level == 0 || (level > 0 && parLevel < OMP_ACTIVE_PARALLEL_LEVEL &&
-                     level <= parLevel)) {
+  if (level == 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL) {
+    rc = threadsInTeam;
+  } else if (level == 0 ||
+             (level > 0 && parLevel < OMP_ACTIVE_PARALLEL_LEVEL &&
+              level <= parLevel) ||
+             (level > 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL &&
+              level <= (parLevel - OMP_ACTIVE_PARALLEL_LEVEL))) {
     rc = 1;
-  } else if (level > 0) {
-    int totLevel = omp_get_level();
-    if (level <= totLevel) {
-      omptarget_nvptx_TaskDescr *currTaskDescr =
-          getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);
-      int steps = totLevel - level;
-      ASSERT0(LT_FUSSY, currTaskDescr,
-              "do not expect fct to be called in a non-active thread");
-      do {
-        if (currTaskDescr->IsParallelConstruct()) {
-          if (!steps) {
-            // found the level
-            rc = currTaskDescr->ThreadsInTeam();
-            break;
-          }
-          steps--;
-        }
-        currTaskDescr = currTaskDescr->GetPrevTaskDescr();
-      } while (currTaskDescr);
-      ASSERT0(LT_FUSSY, !steps, "expected to find all steps");
-    }
   }
   PRINT(LD_IO, "call omp_get_team_size(level %d) returns %d\n", level, rc)
   return rc;
 }
 
 EXTERN void omp_get_schedule(omp_sched_t *kind, int *modifier) {
   if (isRuntimeUninitialized()) {
     ASSERT0(LT_FUSSY, isSPMDMode(),

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

public:
 
   INLINE static void for_static_init(int32_t gtid, int32_t schedtype,
                                      int32_t *plastiter, T *plower, T *pupper,
                                      ST *pstride, ST chunk,
                                      bool IsSPMDExecutionMode) {
     // When IsRuntimeUninitialized is true, we assume that the caller is
     // in an L0 parallel region and that all worker threads participate.
 
-    int tid = GetLogicalThreadIdInBlock(IsSPMDExecutionMode);
-
     // Assume we are in teams region or that we use a single block
     // per target region
-    ST numberOfActiveOMPThreads =
-        GetNumberOfOmpThreads(tid, IsSPMDExecutionMode);
+    ST numberOfActiveOMPThreads = GetNumberOfOmpThreads(IsSPMDExecutionMode);
 
     // All warps that are in excess of the maximum requested, do
     // not execute the loop
     PRINT(LD_LOOP,
           "OMP Thread %d: schedule type %d, chunk size = %lld, mytid "
           "%d, num tids %d\n",
           (int)gtid, (int)schedtype, (long long)chunk, (int)gtid,
           (int)numberOfActiveOMPThreads);
INLINE static void dispatch_init(kmp_Ident *loc, int32_t threadId,
     if (checkRuntimeUninitialized(loc)) {
       // In SPMD mode no need to check parallelism level - dynamic scheduling
       // may appear only in L2 parallel regions with lightweight runtime.
       ASSERT0(LT_FUSSY, checkSPMDMode(loc), "Expected non-SPMD mode.");
       return;
     }
     int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(tid);
-    T tnum = currTaskDescr->ThreadsInTeam();
+    T tnum = GetNumberOfOmpThreads(checkSPMDMode(loc));
     T tripCount = ub - lb + 1; // +1 because ub is inclusive
     ASSERT0(LT_FUSSY, threadId < tnum,
             "current thread is not needed here; error");
 
     /* Currently just ignore the monotonic and non-monotonic modifiers
      * (the compiler isn't producing them * yet anyway).
      * When it is we'll want to look at them somewhere here and use that
      * information to add to our schedule choice. We shouldn't need to pass
if (checkRuntimeUninitialized(loc)) {
         return DISPATCH_FINISHED;
       *plast = 1;
       return DISPATCH_NOTFINISHED;
     }
     // ID of a thread in its own warp
 
     // automatically selects thread or warp ID based on selected implementation
     int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
-    ASSERT0(LT_FUSSY, gtid < GetNumberOfOmpThreads(tid, checkSPMDMode(loc)),
+    ASSERT0(LT_FUSSY, gtid < GetNumberOfOmpThreads(checkSPMDMode(loc)),
             "current thread is not needed here; error");
     // retrieve schedule
     kmp_sched_t schedule =
         omptarget_nvptx_threadPrivateContext->ScheduleType(tid);
 
     // xxx reduce to one
     if (schedule == kmp_sched_static_chunk ||
         schedule == kmp_sched_static_nochunk) {
INLINE static int dispatch_next(kmp_Ident *loc, int32_t gtid, int32_t *plast,
     *plast = (int32_t)(finished == LAST_CHUNK);
     *plower = myLb;
     *pupper = myUb;
     *pstride = 1;
 
     PRINT(LD_LOOP,
           "Got sched: active %d, total %d: lb %lld, ub %lld, stride = %lld, "
           "last %d\n",
-          (int)GetNumberOfOmpThreads(tid, isSPMDMode()),
+          (int)GetNumberOfOmpThreads(isSPMDMode()),
           (int)GetNumberOfWorkersInTeam(), (long long)*plower,
           (long long)*pupper, (long long)*pstride, (int)*plast);
     return DISPATCH_NOTFINISHED;
   }
 
   INLINE static void dispatch_fini() {
     // nothing
   }

 
 EXTERN void __kmpc_reduce_conditional_lastprivate(kmp_Ident *loc, int32_t gtid,
                                                   int32_t varNum, void *array) {
   PRINT0(LD_IO, "call to __kmpc_reduce_conditional_lastprivate(...)\n");
   ASSERT0(LT_FUSSY, checkRuntimeInitialized(loc),
           "Expected non-SPMD mode + initialized runtime.");
 
   omptarget_nvptx_TeamDescr &teamDescr = getMyTeamDescriptor();
-  int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
-  uint32_t NumThreads = GetNumberOfOmpThreads(tid, checkSPMDMode(loc));
+  uint32_t NumThreads = GetNumberOfOmpThreads(checkSPMDMode(loc));
   uint64_t *Buffer = teamDescr.getLastprivateIterBuffer();
   for (unsigned i = 0; i < varNum; i++) {
     // Reset buffer.
     if (gtid == 0)
       *Buffer = 0; // Reset to minimum loop iteration value.
 
     // Barrier.
     syncWorkersInGenericMode(NumThreads);

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

 
 __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 __device__ __shared__ uint32_t usedMemIdx;
 __device__ __shared__ uint32_t usedSlotIdx;
 
 __device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 __device__ __shared__ uint16_t threadLimit;
+__device__ __shared__ uint16_t threadsInTeam;
 // Pointer to this team's OpenMP state object
 __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
 ////////////////////////////////////////////////////////////////////////////////
 // The team master sets the outlined parallel function in this variable to
 // communicate with the workers.  Since it is in shared memory, there is one
 // copy of these variables for each kernel, instance, and team.

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

public:
   }
   INLINE int IsParallelConstruct() const {
     return items.flags & TaskDescr_IsParConstr;
   }
   INLINE int IsTaskConstruct() const { return !IsParallelConstruct(); }
   // methods for other fields
   INLINE uint16_t &NThreads() { return items.nthreads; }
   INLINE uint16_t &ThreadId() { return items.threadId; }
-  INLINE uint16_t &ThreadsInTeam() { return items.threadsInTeam; }
   INLINE uint64_t &RuntimeChunkSize() { return items.runtimeChunkSize; }
   INLINE omptarget_nvptx_TaskDescr *GetPrevTaskDescr() const { return prev; }
   INLINE void SetPrevTaskDescr(omptarget_nvptx_TaskDescr *taskDescr) {
     prev = taskDescr;
   }
   // init & copy
   INLINE void InitLevelZeroTaskDescr(bool isSPMDExecutionMode);
-  INLINE void InitLevelOneTaskDescr(uint16_t tnum,
-                                    omptarget_nvptx_TaskDescr *parentTaskDescr);
+  INLINE void InitLevelOneTaskDescr(omptarget_nvptx_TaskDescr *parentTaskDescr);
   INLINE void Copy(omptarget_nvptx_TaskDescr *sourceTaskDescr);
   INLINE void CopyData(omptarget_nvptx_TaskDescr *sourceTaskDescr);
   INLINE void CopyParent(omptarget_nvptx_TaskDescr *parentTaskDescr);
   INLINE void CopyForExplicitTask(omptarget_nvptx_TaskDescr *parentTaskDescr);
-  INLINE void CopyToWorkDescr(omptarget_nvptx_TaskDescr *masterTaskDescr,
-                              uint16_t tnum);
+  INLINE void CopyToWorkDescr(omptarget_nvptx_TaskDescr *masterTaskDescr);
   INLINE void CopyFromWorkDescr(omptarget_nvptx_TaskDescr *workTaskDescr);
   INLINE void CopyConvergentParent(omptarget_nvptx_TaskDescr *parentTaskDescr,
                                    uint16_t tid, uint16_t tnum);
   INLINE void SaveLoopData();
   INLINE void RestoreLoopData() const;
 
 private:
   // bits for flags: (6 used, 2 free)
struct SavedLoopDescr_items {
     kmp_sched_t schedule;
   } loopData;
 
   struct TaskDescr_items {
     uint8_t flags; // 6 bit used (see flag above)
     uint8_t unused;
     uint16_t nthreads;         // thread num for subsequent parallel regions
     uint16_t threadId;         // thread id
-    uint16_t threadsInTeam;    // threads in current team
     uint64_t runtimeChunkSize; // runtime chunk size
   } items;
   omptarget_nvptx_TaskDescr *prev;
 };
 
 // build on kmp
 typedef struct omptarget_nvptx_ExplicitTaskDescr {
   omptarget_nvptx_TaskDescr

 
 extern __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 extern __device__ __shared__ uint32_t usedMemIdx;
 extern __device__ __shared__ uint32_t usedSlotIdx;
 extern __device__ __shared__ uint8_t
     parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 extern __device__ __shared__ uint16_t threadLimit;
+extern __device__ __shared__ uint16_t threadsInTeam;
 extern __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
 extern __device__ __shared__ uint32_t execution_param;
 extern __device__ __shared__ void *ReductionScratchpadPtr;
 
 ////////////////////////////////////////////////////////////////////////////////
 // work function (outlined parallel/simd functions) and arguments.

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
 
   //
   // Initialize task descr for each thread.
   //
   omptarget_nvptx_TaskDescr *newTaskDescr =
       omptarget_nvptx_threadPrivateContext->Level1TaskDescr(threadId);
   ASSERT0(LT_FUSSY, newTaskDescr, "expected a task descr");
-  newTaskDescr->InitLevelOneTaskDescr(ThreadLimit,
-                                      currTeamDescr.LevelZeroTaskDescr());
+  newTaskDescr->InitLevelOneTaskDescr(currTeamDescr.LevelZeroTaskDescr());
   // install new top descriptor
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                              newTaskDescr);
 
   // init thread private from init value
   PRINT(LD_PAR,
         "thread will execute parallel region with id %d in a team of "
         "%d threads\n",
-        (int)newTaskDescr->ThreadId(), (int)newTaskDescr->ThreadsInTeam());
+        (int)newTaskDescr->ThreadId(), (int)ThreadLimit);
 
   if (RequiresDataSharing && GetLaneId() == 0) {
     // Warp master innitializes data sharing environment.
     unsigned WID = threadId / WARPSIZE;
     __kmpc_data_sharing_slot *RootS = currTeamDescr.RootS(
         WID, WID == WARPSIZE - 1);
     DataSharingState.SlotPtr[WID] = RootS;
     DataSharingState.StackPtr[WID] = (void *)&RootS->Data[0];

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptxi.h

omptarget_nvptx_TaskDescr::InitLevelZeroTaskDescr(bool isSPMDExecutionMode) {
   //   default sched is static,
   //   dyn is off (unused now anyway, but may need to sample from host ?)
   //   not in parallel
 
   items.flags = 0;
   items.nthreads = GetNumberOfProcsInTeam(isSPMDExecutionMode);
   ;                                // threads: whatever was alloc by kernel
   items.threadId = 0;         // is master
-  items.threadsInTeam = 1;    // sequential
   items.runtimeChunkSize = 1; // prefered chunking statik with chunk 1
 }
 
 // This is called when all threads are started together in SPMD mode.
 // OMP directives include target parallel, target distribute parallel for, etc.
 INLINE void omptarget_nvptx_TaskDescr::InitLevelOneTaskDescr(
-    uint16_t tnum, omptarget_nvptx_TaskDescr *parentTaskDescr) {
+    omptarget_nvptx_TaskDescr *parentTaskDescr) {
   // slow method
   // flag:
   //   default sched is static,
   //   dyn is off (unused now anyway, but may need to sample from host ?)
   //   in L1 parallel
 
   items.flags =
       TaskDescr_InPar | TaskDescr_IsParConstr; // set flag to parallel
   items.nthreads = 0; // # threads for subsequent parallel region
   items.threadId =
       GetThreadIdInBlock(); // get ids from cuda (only called for 1st level)
-  items.threadsInTeam = tnum;
   items.runtimeChunkSize = 1; // prefered chunking statik with chunk 1
   prev = parentTaskDescr;
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyData(
     omptarget_nvptx_TaskDescr *sourceTaskDescr) {
   items = sourceTaskDescr->items;
 }

 INLINE void omptarget_nvptx_TaskDescr::CopyForExplicitTask(
     omptarget_nvptx_TaskDescr *parentTaskDescr) {
   CopyParent(parentTaskDescr);
   items.flags = items.flags & ~TaskDescr_IsParConstr;
   ASSERT0(LT_FUSSY, IsTaskConstruct(), "expected task");
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyToWorkDescr(
-    omptarget_nvptx_TaskDescr *masterTaskDescr, uint16_t tnum) {
+    omptarget_nvptx_TaskDescr *masterTaskDescr) {
   CopyParent(masterTaskDescr);
   // overrwrite specific items;
   items.flags |=
       TaskDescr_InPar | TaskDescr_IsParConstr; // set flag to parallel
-  items.threadsInTeam = tnum;             // set number of threads
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyFromWorkDescr(
     omptarget_nvptx_TaskDescr *workTaskDescr) {
   Copy(workTaskDescr);
   //
   // overrwrite specific items;
   //
   // The threadID should be GetThreadIdInBlock() % GetMasterThreadID().
   // This is so that the serial master (first lane in the master warp)
   // gets a threadId of 0.
   // However, we know that this function is always called in a parallel
   // region where only workers are active.  The serial master thread
   // never enters this region.  When a parallel region is executed serially,
   // the threadId is set to 0 elsewhere and the kmpc_serialized_* functions
   // are called, which never activate this region.
   items.threadId =
       GetThreadIdInBlock(); // get ids from cuda (only called for 1st level)
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyConvergentParent(
     omptarget_nvptx_TaskDescr *parentTaskDescr, uint16_t tid, uint16_t tnum) {
   CopyParent(parentTaskDescr);
   items.flags |= TaskDescr_InParL2P; // In L2+ parallelism
-  items.threadsInTeam = tnum;        // set number of threads
   items.threadId = tid;
 }
 
 INLINE void omptarget_nvptx_TaskDescr::SaveLoopData() {
   loopData.loopUpperBound =
       omptarget_nvptx_threadPrivateContext->LoopUpperBound(items.threadId);
   loopData.nextLowerBound =
       omptarget_nvptx_threadPrivateContext->NextLowerBound(items.threadId);

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

EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
 
   ASSERT(LT_FUSSY, NumThreads > 0, "bad thread request of %d threads",
          (int)NumThreads);
   ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "only team master can create parallel");
 
   // Set number of threads on work descriptor.
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
-  workDescr.WorkTaskDescr()->CopyToWorkDescr(currTaskDescr, NumThreads);
+  workDescr.WorkTaskDescr()->CopyToWorkDescr(currTaskDescr);
+  threadsInTeam = NumThreads;
 }
 
 // All workers call this function.  Deactivate those not needed.
 // Fn - the outlined work function to execute.
 // returns True if this thread is active, else False.
 //
 // Only the worker threads call this routine.
 EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
 
   // Only the worker threads call this routine and the master warp
   // never arrives here.  Therefore, use the nvptx thread id.
   int threadId = GetThreadIdInBlock();
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
   // Set to true for workers participating in the parallel region.
   bool isActive = false;
   // Initialize state for active threads.
-  if (threadId < workDescr.WorkTaskDescr()->ThreadsInTeam()) {
+  if (threadId < threadsInTeam) {
     // init work descriptor from workdesccr
     omptarget_nvptx_TaskDescr *newTaskDescr =
         omptarget_nvptx_threadPrivateContext->Level1TaskDescr(threadId);
     ASSERT0(LT_FUSSY, newTaskDescr, "expected a task descr");
     newTaskDescr->CopyFromWorkDescr(workDescr.WorkTaskDescr());
     // install new top descriptor
     omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                                newTaskDescr);
     // init private from int value
     PRINT(LD_PAR,
           "thread will execute parallel region with id %d in a team of "
           "%d threads\n",
           (int)newTaskDescr->ThreadId(), (int)newTaskDescr->NThreads());
 
     isActive = true;
-    IncParallelLevel(workDescr.WorkTaskDescr()->ThreadsInTeam() != 1);
+    IncParallelLevel(threadsInTeam != 1);
   }
 
   return isActive;
 }
 
 EXTERN void __kmpc_kernel_end_parallel() {
   // pop stack
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_end_parallel\n");
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Expected initialized runtime.");
 
   // Only the worker threads call this routine and the master warp
   // never arrives here.  Therefore, use the nvptx thread id.
   int threadId = GetThreadIdInBlock();
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(
       threadId, currTaskDescr->GetPrevTaskDescr());
 
-  DecParallelLevel(currTaskDescr->ThreadsInTeam() != 1);
+  DecParallelLevel(threadsInTeam != 1);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // support for parallel that goes sequential
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_serialized_parallel\n");
omptarget_nvptx_TaskDescr *newTaskDescr =
       (omptarget_nvptx_TaskDescr *)SafeMalloc(sizeof(omptarget_nvptx_TaskDescr),
                                               "new seq parallel task");
   newTaskDescr->CopyParent(currTaskDescr);
 
   // tweak values for serialized parallel case:
   // - each thread becomes ID 0 in its serialized parallel, and
   // - there is only one thread per team
   newTaskDescr->ThreadId() = 0;
-  newTaskDescr->ThreadsInTeam() = 1;
 
   // set new task descriptor as top
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                              newTaskDescr);
 }
 
 EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
                                            uint32_t global_tid) {

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

 #include <stdio.h>
 
 #include "omptarget-nvptx.h"
 
 // may eventually remove this
 EXTERN
 int32_t __gpu_block_reduce() {
   bool isSPMDExecutionMode = isSPMDMode();
-  int tid = GetLogicalThreadIdInBlock(isSPMDExecutionMode);
-  int nt = GetNumberOfOmpThreads(tid, isSPMDExecutionMode);
+  int nt = GetNumberOfOmpThreads(isSPMDExecutionMode);
   if (nt != blockDim.x)
     return 0;
   unsigned tnum = __ACTIVEMASK();
   if (tnum != (~0x0)) // assume swapSize is 32
     return 0;
   return 1;
 }
 
 EXTERN
 int32_t __kmpc_reduce_gpu(kmp_Ident *loc, int32_t global_tid, int32_t num_vars,
                           size_t reduce_size, void *reduce_data,
                           void *reduce_array_size, kmp_ReductFctPtr *reductFct,
                           kmp_CriticalName *lck) {
   int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
   int numthread;
   if (currTaskDescr->IsParallelConstruct()) {
-    numthread = GetNumberOfOmpThreads(threadId, checkSPMDMode(loc));
+    numthread = GetNumberOfOmpThreads(checkSPMDMode(loc));
   } else {
     numthread = GetNumberOfOmpTeams();
   }
 
   if (numthread == 1)
     return 1;
   if (!__gpu_block_reduce())
     return 2;

 }
 
 INLINE
 static 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) {
   uint32_t BlockThreadId = GetLogicalThreadIdInBlock(isSPMDExecutionMode);
-  uint32_t NumThreads =
-      GetNumberOfOmpThreads(BlockThreadId, isSPMDExecutionMode);
+  uint32_t NumThreads = GetNumberOfOmpThreads(isSPMDExecutionMode);
   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.
static int32_t nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
                                          kmp_CopyToScratchpadFctPtr scratchFct,
                                          kmp_LoadReduceFctPtr ldFct,
                                          bool isSPMDExecutionMode) {
   uint32_t ThreadId = GetLogicalThreadIdInBlock(isSPMDExecutionMode);
   // In non-generic mode all workers participate in the teams reduction.
   // In generic mode only the team master participates in the teams
   // reduction because the workers are waiting for parallel work.
   uint32_t NumThreads =
-      isSPMDExecutionMode
-          ? GetNumberOfOmpThreads(ThreadId, /*isSPMDExecutionMode=*/true)
+      isSPMDExecutionMode ? GetNumberOfOmpThreads(/*isSPMDExecutionMode=*/true)
                           : /*Master thread only*/ 1;
   uint32_t TeamId = GetBlockIdInKernel();
   uint32_t NumTeams = GetNumberOfBlocksInKernel();
   __shared__ volatile bool IsLastTeam;
 
   // Team masters of all teams write to the scratchpad.
   if (ThreadId == 0) {
     unsigned int *timestamp = GetTeamsReductionTimestamp();
     char *scratchpad = GetTeamsReductionScratchpad();
if (checkGenericMode(loc) && GetThreadIdInBlock() != GetMasterThreadID())
     return 0;
 
   uint32_t ThreadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
 
   // In non-generic mode all workers participate in the teams reduction.
   // In generic mode only the team master participates in the teams
   // reduction because the workers are waiting for parallel work.
   uint32_t NumThreads =
-      checkSPMDMode(loc)
-          ? GetNumberOfOmpThreads(ThreadId, /*isSPMDExecutionMode=*/true)
+      checkSPMDMode(loc) ? GetNumberOfOmpThreads(/*isSPMDExecutionMode=*/true)
                          : /*Master thread only*/ 1;
   uint32_t TeamId = GetBlockIdInKernel();
   uint32_t NumTeams = GetNumberOfBlocksInKernel();
   __shared__ unsigned Bound;
   __shared__ unsigned ChunkTeamCount;
 
   // Block progress for teams greater than the current upper
   // limit. We always only allow a number of teams less or equal
   // to the number of slots in the buffer.

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/support.h

 INLINE int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
 INLINE int GetOmpThreadId(int threadId,
                           bool isSPMDExecutionMode);    // omp_thread_num
 INLINE int GetOmpTeamId();                              // omp_team_num
 
 // get OpenMP number of threads and team
-INLINE int GetNumberOfOmpThreads(int threadId,
-                                 bool isSPMDExecutionMode); // omp_num_threads
+INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
 INLINE int GetNumberOfOmpTeams();                           // omp_num_teams
 
 // get OpenMP number of procs
 INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
 INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
 
 // masters
 INLINE int IsTeamMaster(int ompThreadId);

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h

if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
     ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");
     rc = currTaskDescr->ThreadId();
   }
   return rc;
 }
 
-INLINE int GetNumberOfOmpThreads(int threadId, bool isSPMDExecutionMode) {
+INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode) {
   // omp_num_threads
   int rc;
-  if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) {
+  int Level = parallelLevel[GetWarpId()];
+  if (Level != OMP_ACTIVE_PARALLEL_LEVEL + 1) {
     rc = 1;
   } else if (isSPMDExecutionMode) {
     rc = GetNumberOfThreadsInBlock();
   } else {
-    omptarget_nvptx_TaskDescr *currTaskDescr =
-        omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
-    ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");
-    rc = currTaskDescr->ThreadsInTeam();
+    rc = threadsInTeam;
   }
 
   return rc;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Team id linked to OpenMP
 

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

 
 EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) {
   if (checkRuntimeUninitialized(loc_ref)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc_ref),
             "Expected SPMD mode with uninitialized runtime.");
     __kmpc_barrier_simple_spmd(loc_ref, tid);
   } else {
     tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc_ref));
-    omptarget_nvptx_TaskDescr *currTaskDescr =
-        omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(tid);
     int numberOfActiveOMPThreads =
-        GetNumberOfOmpThreads(tid, checkSPMDMode(loc_ref));
+        GetNumberOfOmpThreads(checkSPMDMode(loc_ref));
     if (numberOfActiveOMPThreads > 1) {
       if (checkSPMDMode(loc_ref)) {
         __kmpc_barrier_simple_spmd(loc_ref, tid);
       } else {
         // The #threads parameter must be rounded up to the WARPSIZE.
         int threads =
             WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);