[OPENMP][NVPTX]Improve thread limit counter, NFC.

[ABataev]

Summary

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

[grokos]

Looks good.

libomptarget/deviceRTLs/nvptx/src/libcall.cu

EXTERN int omp_get_max_threads(void) {
   ASSERT0(LT_FUSSY, rc >= 0, "bad number of threads");
   PRINT(LD_IO, "call omp_get_max_threads() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_thread_limit(void) {
   if (isSPMDMode())
     return GetNumberOfThreadsInBlock();
-  // per contention group.. meaning threads in current team
-  omptarget_nvptx_TaskDescr *currTaskDescr =
-      getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);
-  int rc = currTaskDescr->ThreadLimit();
+  int rc = threadLimit;
   PRINT(LD_IO, "call omp_get_thread_limit() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_thread_num() {
   bool isSPMDExecutionMode = isSPMDMode();
   int tid = GetLogicalThreadIdInBlock(isSPMDExecutionMode);
   int rc = GetOmpThreadId(tid, isSPMDExecutionMode);

libomptarget/deviceRTLs/nvptx/src/omp_data.cu

omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
 __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;
 // 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.

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

INLINE int InL2OrHigherParallelRegion() const {
     return items.flags & TaskDescr_InParL2P;
   }
   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 &ThreadLimit() { return items.threadlimit; }
   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
struct SavedLoopDescr_items {
     int64_t stride;
     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 threadlimit;      // thread limit ICV
     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

 ////////////////////////////////////////////////////////////////////////////////
 
 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__
     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.

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

EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime) {
   // InitTeamDescr()
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(
       threadId, currTeamDescr.LevelZeroTaskDescr());
 
   // set number of threads and thread limit in team to started value
   omptarget_nvptx_TaskDescr *currTaskDescr =
       omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
   currTaskDescr->NThreads() = GetNumberOfWorkersInTeam();
-  currTaskDescr->ThreadLimit() = ThreadLimit;
+  threadLimit = ThreadLimit;
 }
 
 EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO, "call to __kmpc_kernel_deinit\n");
   ASSERT0(LT_FUSSY, IsOMPRuntimeInitialized,
           "Generic always requires initialized runtime.");
   // Enqueue omp state object for use by another team.
   int slot = usedSlotIdx;
EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
   //
   // 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->ThreadLimit() = ThreadLimit;
   // 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",

libomptarget/deviceRTLs/nvptx/src/parallel.cu

EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
   if (currTaskDescr->InParallelRegion()) {
     PRINT0(LD_PAR, "already in parallel: go seq\n");
     return;
   }
 
   uint16_t &NumThreadsClause =
       omptarget_nvptx_threadPrivateContext->NumThreadsForNextParallel(threadId);
 
-  uint16_t NumThreads =
-      determineNumberOfThreads(NumThreadsClause, currTaskDescr->NThreads(),
-                               currTaskDescr->ThreadLimit());
+  uint16_t NumThreads = determineNumberOfThreads(
+      NumThreadsClause, currTaskDescr->NThreads(), threadLimit);
 
   if (NumThreadsClause != 0) {
     // Reset request to avoid propagating to successive #parallel
     NumThreadsClause = 0;
   }
 
   ASSERT(LT_FUSSY, NumThreads > 0, "bad thread request of %d threads",
          (int)NumThreads);