[OpenMP][libomptarget] Use shared memory variable for tracking parallel level

[gtbercea]

Summary

Replace existing infrastructure for tracking parallel level using global memory with a per-team shared memory variable. This minimizes the impact of the overhead of tracking the parallel level for non-nested cases.

[gtbercea]

ping

[ABataev]

libomptarget/deviceRTLs/nvptx/src/parallel.cu

345

1. At first, you need to synchronize the threads, and only after that, you can increment the counter.
2. Use __SYNCTHREADS() instead of the __syncthreads().
3. Use preincrement.

386–389

1. Predecrement.
2. Synchronize at first, then decerement.
3. Use __SYNCTHREADS()

[gtbercea]

• Use new sync.

[ABataev]

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

101–102

1. I think it is better to do this initialization in the single thread and you still need to synchronize the threads even in this situation to avoid data race.
2. Keep the initialization of the usedSlotIdx, it is used in other places.

[gtbercea]

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

101–102

Increments/decrements to this variable are guarded by syncthreads so I don't think we need any additional synchronization here. As for the slot I think we can leave that to have the default value for this case and rely on the compiler to perform the optimization of the remainder operations.

[gtbercea]

• Add slot initialization.

[ABataev]

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

98

Better to put this initialization under the guard of the next if statement.

[gtbercea]

• Fix.

[gtbercea]

• Fix.

[ABataev]

LG

libomptarget/deviceRTLs/nvptx/src/libcall.cu

EXTERN int omp_get_max_active_levels(void) {
   PRINT(LD_IO, "call omp_get_max_active_levels() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_level(void) {
   if (isRuntimeUninitialized()) {
     ASSERT0(LT_FUSSY, isSPMDMode(),
             "Expected SPMD mode only with uninitialized runtime.");
-    return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
+    return parallelLevel;
   }
   int level = 0;
   omptarget_nvptx_TaskDescr *currTaskDescr =
       getMyTopTaskDescriptor(isSPMDMode());
   ASSERT0(LT_FUSSY, currTaskDescr,
           "do not expect fct to be called in a non-active thread");
   do {
     if (currTaskDescr->IsParallelConstruct()) {

libomptarget/deviceRTLs/nvptx/src/omp_data.cu

 ////////////////////////////////////////////////////////////////////////////////
 // global data holding OpenMP state information
 ////////////////////////////////////////////////////////////////////////////////
 
 __device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
-__device__ omptarget_nvptx_Queue<omptarget_nvptx_SimpleThreadPrivateContext,
-                                 OMP_STATE_COUNT>
-    omptarget_nvptx_device_simpleState[MAX_SM];
-
 __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 __device__ __shared__ uint32_t usedMemIdx;
 __device__ __shared__ uint32_t usedSlotIdx;
 
+__device__ __shared__ uint8_t parallelLevel;
+
 // Pointer to this team's OpenMP state object
 __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
-__device__ __shared__ omptarget_nvptx_SimpleThreadPrivateContext
-    *omptarget_nvptx_simpleThreadPrivateContext;
-
 ////////////////////////////////////////////////////////////////////////////////
 // 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.
 ////////////////////////////////////////////////////////////////////////////////
 volatile __device__ __shared__ omptarget_nvptx_WorkFn omptarget_nvptx_workFn;
 
 ////////////////////////////////////////////////////////////////////////////////

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

INLINE static uint32_t hash(unsigned key) {
     return key & (OMP_STATE_COUNT - 1);
   }
 
 public:
   INLINE void Release();
   INLINE const void *Acquire(const void *buf, size_t size);
 };
 
-class omptarget_nvptx_SimpleThreadPrivateContext {
-  uint16_t par_level[MAX_THREADS_PER_TEAM];
-
-public:
-  INLINE void Init() {
-    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
-            "Expected SPMD + uninitialized runtime modes.");
-    par_level[GetThreadIdInBlock()] = 0;
-  }
-  INLINE void IncParLevel() {
-    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
-            "Expected SPMD + uninitialized runtime modes.");
-    ++par_level[GetThreadIdInBlock()];
-  }
-  INLINE void DecParLevel() {
-    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
-            "Expected SPMD + uninitialized runtime modes.");
-    ASSERT0(LT_FUSSY, par_level[GetThreadIdInBlock()] > 0,
-            "Expected parallel level >0.");
-    --par_level[GetThreadIdInBlock()];
-  }
-  INLINE bool InL2OrHigherParallelRegion() const {
-    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
-            "Expected SPMD + uninitialized runtime modes.");
-    return par_level[GetThreadIdInBlock()] > 0;
-  }
-  INLINE uint16_t GetParallelLevel() const {
-    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
-            "Expected SPMD + uninitialized runtime modes.");
-    return par_level[GetThreadIdInBlock()] + 1;
-  }
-};
-
 ////////////////////////////////////////////////////////////////////////////////
 // global device envrionment
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__ omptarget_device_environmentTy omptarget_device_environment;
 
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 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;
 extern __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
-extern __device__ __shared__ omptarget_nvptx_SimpleThreadPrivateContext
-    *omptarget_nvptx_simpleThreadPrivateContext;
 
 extern __device__ __shared__ uint32_t execution_param;
 extern __device__ __shared__ void *ReductionScratchpadPtr;
 
 ////////////////////////////////////////////////////////////////////////////////
 // work function (outlined parallel/simd functions) and arguments.
 // needed for L1 parallelism only.
 ////////////////////////////////////////////////////////////////////////////////

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

 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
-extern __device__ omptarget_nvptx_Queue<
-    omptarget_nvptx_SimpleThreadPrivateContext, OMP_STATE_COUNT>
-    omptarget_nvptx_device_simpleState[MAX_SM];
-
 ////////////////////////////////////////////////////////////////////////////////
 // init entry points
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE static unsigned smid() {
   unsigned id;
   asm("mov.u32 %0, %%smid;" : "=r"(id));
   return id;

 
 EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
                                     int16_t RequiresDataSharing) {
   PRINT0(LD_IO, "call to __kmpc_spmd_kernel_init\n");
 
   if (!RequiresOMPRuntime) {
     // If OMP runtime is not required don't initialize OMP state.
     setExecutionParameters(Spmd, RuntimeUninitialized);
-    if (GetThreadIdInBlock() == 0) {
-      int slot = smid() % MAX_SM;
-      usedSlotIdx = slot;
-      omptarget_nvptx_simpleThreadPrivateContext =
-          omptarget_nvptx_device_simpleState[slot].Dequeue();
-    }
-    // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+    parallelLevel = 0;
+    if (threadId == 0)
+      usedSlotIdx = smid() % MAX_SM;
     __SYNCTHREADS();
-    omptarget_nvptx_simpleThreadPrivateContext->Init();
     return;
   }
   setExecutionParameters(Spmd, RuntimeInitialized);
 
   //
   // Team Context Initialization.
   //
   // In SPMD mode there is no master thread so use any cuda thread for team
   // context initialization.

 
 EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit() {
   __kmpc_spmd_kernel_deinit_v2(isRuntimeInitialized());
 }
 
 EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime) {
   // We're not going to pop the task descr stack of each thread since
   // there are no more parallel regions in SPMD mode.
+  if (!RequiresOMPRuntime)
+    return;
+
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
   __SYNCTHREADS();
   int threadId = GetThreadIdInBlock();
-  if (!RequiresOMPRuntime) {
-    if (threadId == 0) {
-      // Enqueue omp state object for use by another team.
-      int slot = usedSlotIdx;
-      omptarget_nvptx_device_simpleState[slot].Enqueue(
-          omptarget_nvptx_simpleThreadPrivateContext);
-    }
-    return;
-  }
   if (threadId == 0) {
     // Enqueue omp state object for use by another team.
     int slot = usedSlotIdx;
     omptarget_nvptx_device_State[slot].Enqueue(
         omptarget_nvptx_threadPrivateContext);
   }
 }
 
 // Return true if the current target region is executed in SPMD mode.
 EXTERN int8_t __kmpc_is_spmd_exec_mode() {
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_is_spmd_exec_mode\n");
   return isSPMDMode();
 }

libomptarget/deviceRTLs/nvptx/src/parallel.cu

 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_serialized_parallel\n");
 
   if (checkRuntimeUninitialized(loc)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc),
             "Expected SPMD mode with uninitialized runtime.");
-    omptarget_nvptx_simpleThreadPrivateContext->IncParLevel();
+    __SYNCTHREADS();
+    if (GetThreadIdInBlock() == 0)
+      ++parallelLevel;
+    __SYNCTHREADS();
+
     return;
   }
 
   // assume this is only called for nested parallel
   int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
 
   // unlike actual parallel, threads in the same team do not share
   // the workTaskDescr in this case and num threads is fixed to 1

 
 EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
                                            uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_end_serialized_parallel\n");
 
   if (checkRuntimeUninitialized(loc)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc),
             "Expected SPMD mode with uninitialized runtime.");
-    omptarget_nvptx_simpleThreadPrivateContext->DecParLevel();
+    __SYNCTHREADS();
+    if (GetThreadIdInBlock() == 0)
+      --parallelLevel;
+    __SYNCTHREADS();
     return;
   }
 
   // pop stack
   int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
   // set new top
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(
       threadId, currTaskDescr->GetPrevTaskDescr());
   // free
   SafeFree(currTaskDescr, (char *)"new seq parallel task");
   currTaskDescr = getMyTopTaskDescriptor(threadId);
   currTaskDescr->RestoreLoopData();
 }
 
 EXTERN uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_parallel_level\n");
 
   if (checkRuntimeUninitialized(loc)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc),
             "Expected SPMD mode with uninitialized runtime.");
-    return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
+    return parallelLevel;
   }
 
   int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
   omptarget_nvptx_TaskDescr *currTaskDescr =
       omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
   if (currTaskDescr->InL2OrHigherParallelRegion())
     return 2;
   else if (currTaskDescr->InParallelRegion())

libomptarget/deviceRTLs/nvptx/src/supporti.h

INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode,
                           bool isRuntimeUninitialized) {
   // omp_thread_num
   int rc;
 
   if (isRuntimeUninitialized) {
     ASSERT0(LT_FUSSY, isSPMDExecutionMode,
             "Uninitialized runtime with non-SPMD mode.");
     // For level 2 parallelism all parallel regions are executed sequentially.
-    if (omptarget_nvptx_simpleThreadPrivateContext
-            ->InL2OrHigherParallelRegion())
+    if (parallelLevel > 0)
       rc = 0;
     else
       rc = GetThreadIdInBlock();
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
     rc = currTaskDescr->ThreadId();
   }
   return rc;
 }
 
 INLINE int GetNumberOfOmpThreads(int threadId, bool isSPMDExecutionMode,
                                  bool isRuntimeUninitialized) {
   // omp_num_threads
   int rc;
 
   if (isRuntimeUninitialized) {
     ASSERT0(LT_FUSSY, isSPMDExecutionMode,
             "Uninitialized runtime with non-SPMD mode.");
     // For level 2 parallelism all parallel regions are executed sequentially.
-    if (omptarget_nvptx_simpleThreadPrivateContext
-            ->InL2OrHigherParallelRegion())
+    if (parallelLevel > 0)
       rc = 1;
     else
       rc = GetNumberOfThreadsInBlock();
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
     ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");
     rc = currTaskDescr->ThreadsInTeam();