[OPENMP][NVPTX]Improve code by using parallel level counter.

[ABataev]

Summary

Previously for the different purposes we need to get the active/common
parallel level and with full runtime we iterated over all the records to
calculate this level. Instead, we can used the warp-based parallel level
counters used in no-runtime mode.

[ABataev]

Fixed detection of non-active parallel regions.

[ABataev]

Update to the latest version

[grokos]

LGTM

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

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, isRuntimeUninitialized());
+  int rc = GetNumberOfOmpThreads(tid, isSPMDExecutionMode);
   PRINT(LD_IO, "call omp_get_num_threads() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_max_threads(void) {
   if (isSPMDMode())
     // We're already in parallel region.
     return 1;  // default is 1 thread avail
EXTERN int omp_get_thread_limit(void) {
   int rc = currTaskDescr->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, isRuntimeUninitialized());
+  int rc = GetOmpThreadId(tid, isSPMDExecutionMode);
   PRINT(LD_IO, "call omp_get_thread_num() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_num_procs(void) {
   int rc = GetNumberOfProcsInDevice(isSPMDMode());
   PRINT(LD_IO, "call omp_get_num_procs() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_in_parallel(void) {
-  int rc = 0;
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
-    rc = 1;  // SPMD mode is always in parallel.
-  } else {
-    omptarget_nvptx_TaskDescr *currTaskDescr =
-        getMyTopTaskDescriptor(isSPMDMode());
-    if (currTaskDescr->InParallelRegion()) {
-      rc = 1;
-    }
-  }
+  int rc = parallelLevel[GetWarpId()] > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0;
   PRINT(LD_IO, "call omp_in_parallel() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_in_final(void) {
   // treat all tasks as final... Specs may expect runtime to keep
   // track more precisely if a task was actively set by users... This
   // is not explicitely specified; will treat as if runtime can

 
 EXTERN int omp_get_max_active_levels(void) {
   int rc = 1;
   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.");
-    // parallelLevel starts from 0, need to add 1 for correct level.
-    return parallelLevel[GetWarpId()] + 1;
-  }
-  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()) {
-      level++;
-    }
-    currTaskDescr = currTaskDescr->GetPrevTaskDescr();
-  } while (currTaskDescr);
+  int level = parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1);
   PRINT(LD_IO, "call omp_get_level() returns %d\n", level);
   return level;
 }
 
 EXTERN int omp_get_active_level(void) {
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
-    return 1;
-  }
-  int level = 0; // no active level parallelism
-  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->ThreadsInTeam() > 1) {
-      // has a parallel with more than one thread in team
-      level = 1;
-      break;
-    }
-    currTaskDescr = currTaskDescr->GetPrevTaskDescr();
-  } while (currTaskDescr);
+  int level = parallelLevel[GetWarpId()] > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0;
   PRINT(LD_IO, "call omp_get_active_level() returns %d\n", level)
   return level;
 }
 
 EXTERN int omp_get_ancestor_thread_num(int level) {
   if (isRuntimeUninitialized()) {
     ASSERT0(LT_FUSSY, isSPMDMode(),
             "Expected SPMD mode only with uninitialized runtime.");

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

public:
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for Static Init
 
   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,
-                                     bool IsRuntimeUninitialized) {
+                                     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, IsRuntimeUninitialized);
+    ST numberOfActiveOMPThreads =
+        GetNumberOfOmpThreads(tid, 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);
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),
-                                         checkRuntimeUninitialized(loc)),
+    ASSERT0(LT_FUSSY, gtid < GetNumberOfOmpThreads(tid, 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) {
if (finished == FINISHED)
       return DISPATCH_FINISHED;
 
     // not finished (either not finished or last chunk)
     *plast = (int32_t)(finished == LAST_CHUNK);
     *plower = myLb;
     *pupper = myUb;
     *pstride = 1;
 
-    PRINT(
-        LD_LOOP,
+    PRINT(LD_LOOP,
           "Got sched: active %d, total %d: lb %lld, ub %lld, stride = %lld, "
           "last %d\n",
-        (int)GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
-        (int)GetNumberOfWorkersInTeam(), (long long)*plower, (long long)*pupper,
-        (long long)*pstride, (int)*plast);
+          (int)GetNumberOfOmpThreads(tid, 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_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
                                      int32_t schedtype, int32_t *plastiter,
                                      int32_t *plower, int32_t *pupper,
                                      int32_t *pstride, int32_t incr,
                                      int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      checkSPMDMode(loc), checkRuntimeUninitialized(loc));
+      checkSPMDMode(loc));
 }
 
 EXTERN void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
                                       int32_t schedtype, int32_t *plastiter,
                                       uint32_t *plower, uint32_t *pupper,
                                       int32_t *pstride, int32_t incr,
                                       int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      checkSPMDMode(loc), checkRuntimeUninitialized(loc));
+      checkSPMDMode(loc));
 }
 
 EXTERN void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
                                      int32_t schedtype, int32_t *plastiter,
                                      int64_t *plower, int64_t *pupper,
                                      int64_t *pstride, int64_t incr,
                                      int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      checkSPMDMode(loc), checkRuntimeUninitialized(loc));
+      checkSPMDMode(loc));
 }
 
 EXTERN void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
                                       int32_t schedtype, int32_t *plastiter,
                                       uint64_t *plower, uint64_t *pupper,
                                       int64_t *pstride, int64_t incr,
                                       int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      checkSPMDMode(loc), checkRuntimeUninitialized(loc));
+      checkSPMDMode(loc));
 }
 
 EXTERN
 void __kmpc_for_static_init_4_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t schedtype, int32_t *plastiter,
                                           int32_t *plower, int32_t *pupper,
                                           int32_t *pstride, int32_t incr,
                                           int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4_simple_spmd\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/true, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/true);
 }
 
 EXTERN
 void __kmpc_for_static_init_4u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t schedtype,
                                            int32_t *plastiter, uint32_t *plower,
                                            uint32_t *pupper, int32_t *pstride,
                                            int32_t incr, int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u_simple_spmd\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/true, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/true);
 }
 
 EXTERN
 void __kmpc_for_static_init_8_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t schedtype, int32_t *plastiter,
                                           int64_t *plower, int64_t *pupper,
                                           int64_t *pstride, int64_t incr,
                                           int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8_simple_spmd\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/true, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/true);
 }
 
 EXTERN
 void __kmpc_for_static_init_8u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t schedtype,
                                            int32_t *plastiter, uint64_t *plower,
                                            uint64_t *pupper, int64_t *pstride,
                                            int64_t incr, int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u_simple_spmd\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/true, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/true);
 }
 
 EXTERN
 void __kmpc_for_static_init_4_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     int32_t *plower, int32_t *pupper, int32_t *pstride, int32_t incr,
     int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4_simple_generic\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/false, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/false);
 }
 
 EXTERN
 void __kmpc_for_static_init_4u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     uint32_t *plower, uint32_t *pupper, int32_t *pstride, int32_t incr,
     int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u_simple_generic\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/false, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/false);
 }
 
 EXTERN
 void __kmpc_for_static_init_8_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     int64_t *plower, int64_t *pupper, int64_t *pstride, int64_t incr,
     int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8_simple_generic\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/false, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/false);
 }
 
 EXTERN
 void __kmpc_for_static_init_8u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     uint64_t *plower, uint64_t *pupper, int64_t *pstride, int64_t incr,
     int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u_simple_generic\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
-      /*IsSPMDExecutionMode=*/false, /*IsRuntimeUninitialized=*/true);
+      /*IsSPMDExecutionMode=*/false);
 }
 
 EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid) {
   PRINT0(LD_IO, "call kmpc_for_static_fini\n");
 }
 
 namespace {
 INLINE void syncWorkersInGenericMode(uint32_t NumThreads) {

 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),
-                                              checkRuntimeUninitialized(loc));
+  uint32_t NumThreads = GetNumberOfOmpThreads(tid, 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/omptarget-nvptx.cu

 }
 
 EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime) {
   PRINT(LD_IO, "call to __kmpc_kernel_init with version %f\n",
         OMPTARGET_NVPTX_VERSION);
   ASSERT0(LT_FUSSY, RequiresOMPRuntime,
           "Generic always requires initialized runtime.");
   setExecutionParameters(Generic, RuntimeInitialized);
+  for (int I = 0; I < MAX_THREADS_PER_TEAM / WARPSIZE; ++I)
+    parallelLevel[I] = 0;
 
   int threadIdInBlock = GetThreadIdInBlock();
   ASSERT0(LT_FUSSY, threadIdInBlock == GetMasterThreadID(),
           "__kmpc_kernel_init() must be called by team master warp only!");
   PRINT0(LD_IO, "call to __kmpc_kernel_init for master\n");
 
   // Get a state object from the queue.
   int slot = smid() % MAX_SM;
EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized) {
   // Done with work.  Kill the workers.
   omptarget_nvptx_workFn = 0;
 }
 
 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) {
-      usedSlotIdx = smid() % MAX_SM;
-      parallelLevel[0] = 0;
-    } else if (GetLaneId() == 0) {
-      parallelLevel[GetWarpId()] = 0;
-    }
+  setExecutionParameters(Spmd, RequiresOMPRuntime ? RuntimeInitialized
+                                                  : RuntimeUninitialized);
+  int threadId = GetThreadIdInBlock();
+  if (threadId == 0) {
+    parallelLevel[0] =
+        1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
+    parallelLevel[GetWarpId()] =
+        1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
+  if (!RequiresOMPRuntime) {
+    // Runtime is not required - exit.
     __SYNCTHREADS();
     return;
   }
-  setExecutionParameters(Spmd, RuntimeInitialized);
 
   //
   // Team Context Initialization.
   //
   // In SPMD mode there is no master thread so use any cuda thread for team
   // context initialization.
-  int threadId = GetThreadIdInBlock();
   if (threadId == 0) {
     // Get a state object from the queue.
-    int slot = smid() % MAX_SM;
-    usedSlotIdx = slot;
     omptarget_nvptx_threadPrivateContext =
-        omptarget_nvptx_device_State[slot].Dequeue();
+        omptarget_nvptx_device_State[usedSlotIdx].Dequeue();
 
     omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
     omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
     // init team context
     currTeamDescr.InitTeamDescr(/*isSPMDExecutionMode=*/true);
   }
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
   __SYNCTHREADS();
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());
 
-  if (RequiresDataSharing && threadId % WARPSIZE == 0) {
+  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/option.h

 #elif __CUDA_ARCH__ >= 600
 #define OMP_STATE_COUNT 32
 #define MAX_SM 56
 #else
 #define OMP_STATE_COUNT 16
 #define MAX_SM 16
 #endif
 
+#define OMP_ACTIVE_PARALLEL_LEVEL 128
+
 ////////////////////////////////////////////////////////////////////////////////
 // algo options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // misc options (by def everythig here is device)
 ////////////////////////////////////////////////////////////////////////////////
 

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

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);
   }
 
   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);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // 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");
 
+  IncParallelLevel(/*ActiveParallel=*/false);
+
   if (checkRuntimeUninitialized(loc)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc),
             "Expected SPMD mode with uninitialized runtime.");
-    unsigned tnum = __ACTIVEMASK();
-    int leader = __ffs(tnum) - 1;
-    __SHFL_SYNC(tnum, leader, leader);
-    if (GetLaneId() == leader)
-      ++parallelLevel[GetWarpId()];
-    __SHFL_SYNC(tnum, leader, leader);
-
     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_serialized_parallel(kmp_Ident *loc, uint32_t global_tid) {
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                              newTaskDescr);
 }
 
 EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
                                            uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_end_serialized_parallel\n");
 
+  DecParallelLevel(/*ActiveParallel=*/false);
+
   if (checkRuntimeUninitialized(loc)) {
     ASSERT0(LT_FUSSY, checkSPMDMode(loc),
             "Expected SPMD mode with uninitialized runtime.");
-    unsigned tnum = __ACTIVEMASK();
-    int leader = __ffs(tnum) - 1;
-    __SHFL_SYNC(tnum, leader, leader);
-    if (GetLaneId() == leader)
-      --parallelLevel[GetWarpId()];
-    __SHFL_SYNC(tnum, leader, leader);
     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 parallelLevel[GetWarpId()] + 1;
-  }
-
-  int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
-  omptarget_nvptx_TaskDescr *currTaskDescr =
-      omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
-  if (currTaskDescr->InL2OrHigherParallelRegion())
-    return 2;
-  else if (currTaskDescr->InParallelRegion())
-    return 1;
-  else
-    return 0;
+  return parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1);
 }
 
 // This kmpc call returns the thread id across all teams. It's value is
 // cached by the compiler and used when calling the runtime. On nvptx
 // it's cheap to recalculate this value so we never use the result
 // of this call.
 EXTERN int32_t __kmpc_global_thread_num(kmp_Ident *loc) {
   int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
-  return GetOmpThreadId(tid, checkSPMDMode(loc),
-                        checkRuntimeUninitialized(loc));
+  return GetOmpThreadId(tid, checkSPMDMode(loc));
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // push params
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_push_num_threads(kmp_Ident *loc, int32_t tid,
                                     int32_t num_threads) {

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

 
 #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, isRuntimeUninitialized());
+  int nt = GetNumberOfOmpThreads(tid, 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),
-                              checkRuntimeUninitialized(loc));
+    numthread = GetNumberOfOmpThreads(threadId, 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, isRuntimeUninitialized);
+  uint32_t NumThreads =
+      GetNumberOfOmpThreads(BlockThreadId, 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.
#else
   // an L2 parallel region.
   return global_tid == 0;
 #endif // __CUDA_ARCH__ >= 700
 }
 
 EXTERN __attribute__((deprecated)) int32_t __kmpc_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) {
-  return nvptx_parallel_reduce_nowait(
-      global_tid, num_vars, reduce_size, reduce_data, shflFct, cpyFct,
-      /*isSPMDExecutionMode=*/isSPMDMode(),
-      /*isRuntimeUninitialized=*/isRuntimeUninitialized());
+  return nvptx_parallel_reduce_nowait(global_tid, num_vars, reduce_size,
+                                      reduce_data, shflFct, cpyFct,
+                                      isSPMDMode(), isRuntimeUninitialized());
 }
 
 EXTERN
 int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
     kmp_Ident *loc, int32_t global_tid, int32_t num_vars, size_t reduce_size,
     void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
     kmp_InterWarpCopyFctPtr cpyFct) {
   return nvptx_parallel_reduce_nowait(
       global_tid, num_vars, reduce_size, reduce_data, shflFct, cpyFct,
       checkSPMDMode(loc), checkRuntimeUninitialized(loc));
 }
 
 EXTERN
 int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_spmd(
     int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
     kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct) {
-  return nvptx_parallel_reduce_nowait(global_tid, num_vars, reduce_size,
-                                      reduce_data, shflFct, cpyFct,
-                                      /*isSPMDExecutionMode=*/true,
-                                      /*isRuntimeUninitialized=*/true);
+  return nvptx_parallel_reduce_nowait(
+      global_tid, num_vars, reduce_size, reduce_data, shflFct, cpyFct,
+      /*isSPMDExecutionMode=*/true, /*isRuntimeUninitialized=*/true);
 }
 
 EXTERN
 int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_generic(
     int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
     kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct) {
-  return nvptx_parallel_reduce_nowait(global_tid, num_vars, reduce_size,
-                                      reduce_data, shflFct, cpyFct,
-                                      /*isSPMDExecutionMode=*/false,
-                                      /*isRuntimeUninitialized=*/true);
+  return nvptx_parallel_reduce_nowait(
+      global_tid, num_vars, reduce_size, reduce_data, shflFct, cpyFct,
+      /*isSPMDExecutionMode=*/false, /*isRuntimeUninitialized=*/true);
 }
 
 INLINE
-static int32_t nvptx_teams_reduce_nowait(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
-    kmp_CopyToScratchpadFctPtr scratchFct, kmp_LoadReduceFctPtr ldFct,
-    bool isSPMDExecutionMode, bool isRuntimeUninitialized) {
+static int32_t nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
+                                         size_t reduce_size, void *reduce_data,
+                                         kmp_ShuffleReductFctPtr shflFct,
+                                         kmp_InterWarpCopyFctPtr cpyFct,
+                                         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,
-                                  isRuntimeUninitialized)
+          ? GetNumberOfOmpThreads(ThreadId, /*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();

 
 EXTERN
 int32_t __kmpc_nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
                                          size_t reduce_size, void *reduce_data,
                                          kmp_ShuffleReductFctPtr shflFct,
                                          kmp_InterWarpCopyFctPtr cpyFct,
                                          kmp_CopyToScratchpadFctPtr scratchFct,
                                          kmp_LoadReduceFctPtr ldFct) {
-  return nvptx_teams_reduce_nowait(
-      global_tid, num_vars, reduce_size, reduce_data, shflFct, cpyFct,
-      scratchFct, ldFct, /*isSPMDExecutionMode=*/isSPMDMode(),
-      /*isRuntimeUninitialized=*/isRuntimeUninitialized());
+  return nvptx_teams_reduce_nowait(global_tid, num_vars, reduce_size,
+                                   reduce_data, shflFct, cpyFct, scratchFct,
+                                   ldFct, isSPMDMode());
 }
 
 EXTERN
 int32_t __kmpc_nvptx_teams_reduce_nowait_simple_spmd(
     int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
     kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
     kmp_CopyToScratchpadFctPtr scratchFct, kmp_LoadReduceFctPtr ldFct) {
   return nvptx_teams_reduce_nowait(global_tid, num_vars, reduce_size,
                                    reduce_data, shflFct, cpyFct, scratchFct,
-                                   ldFct,
-                                   /*isSPMDExecutionMode=*/true,
-                                   /*isRuntimeUninitialized=*/true);
+                                   ldFct, /*isSPMDExecutionMode=*/true);
 }
 
 EXTERN
 int32_t __kmpc_nvptx_teams_reduce_nowait_simple_generic(
     int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
     kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
     kmp_CopyToScratchpadFctPtr scratchFct, kmp_LoadReduceFctPtr ldFct) {
   return nvptx_teams_reduce_nowait(global_tid, num_vars, reduce_size,
                                    reduce_data, shflFct, cpyFct, scratchFct,
-                                   ldFct,
-                                   /*isSPMDExecutionMode=*/false,
-                                   /*isRuntimeUninitialized=*/true);
+                                   ldFct, /*isSPMDExecutionMode=*/false);
 }
 
 EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple(kmp_Ident *loc,
                                                        int32_t global_tid,
                                                        kmp_CriticalName *crit) {
   if (checkSPMDMode(loc) && GetThreadIdInBlock() != 0)
     return 0;
   // The master thread of the team actually does the reduction.
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
 
   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,
-                                  checkRuntimeUninitialized(loc))
+          ? GetNumberOfOmpThreads(ThreadId, /*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

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

 INLINE unsigned GetLaneId();
 
 // get global ids to locate tread/team info (constant regardless of OMP)
 INLINE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
 INLINE int GetMasterThreadID();
 INLINE int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
-INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode,
-                          bool isRuntimeUninitialized); // omp_thread_num
+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,
-                      bool isRuntimeUninitialized); // omp_num_threads
-INLINE int GetNumberOfOmpTeams();                   // omp_num_teams
+INLINE int GetNumberOfOmpThreads(int threadId,
+                                 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);
 
+// Parallel level
+INLINE void IncParallelLevel(bool ActiveParallel);
+INLINE void DecParallelLevel(bool ActiveParallel);
+
 ////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
 // safe alloc and free
 INLINE void *SafeMalloc(size_t size, const char *msg); // check if success
 INLINE void *SafeFree(void *ptr, const char *msg);
 // pad to a alignment (power of 2 only)

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

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // OpenMP Thread Support Layer
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode,
-                          bool isRuntimeUninitialized) {
+INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) {
   // 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 (parallelLevel[GetWarpId()] > 0)
-      rc = 0;
-    else
-      rc = GetThreadIdInBlock();
+  if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) {
+  } else if (isSPMDExecutionMode) {
   } else {
     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,
-                                 bool isRuntimeUninitialized) {
+INLINE int GetNumberOfOmpThreads(int threadId, bool isSPMDExecutionMode) {
   // 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 (parallelLevel[GetWarpId()] > 0)
-      rc = 1;
-    else
-      rc = GetNumberOfThreadsInBlock();
+  if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) {
+  } else if (isSPMDExecutionMode) {
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
     ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");
     rc = currTaskDescr->ThreadsInTeam();
   }
 
   return rc;

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Masters
 
 INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 ////////////////////////////////////////////////////////////////////////////////
+// Parallel level
+
+INLINE void IncParallelLevel(bool ActiveParallel) {
+  unsigned tnum = __ACTIVEMASK();
+  int leader = __ffs(tnum) - 1;
+  __SHFL_SYNC(tnum, leader, leader);
+  if (GetLaneId() == leader) {
+    parallelLevel[GetWarpId()] +=
+        (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
+  }
+  __SHFL_SYNC(tnum, leader, leader);
+}
+
+INLINE void DecParallelLevel(bool ActiveParallel) {
+  unsigned tnum = __ACTIVEMASK();
+  int leader = __ffs(tnum) - 1;
+  __SHFL_SYNC(tnum, leader, leader);
+  if (GetLaneId() == leader) {
+    parallelLevel[GetWarpId()] -=
+        (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
+  }
+  __SHFL_SYNC(tnum, leader, leader);
+}
+
+////////////////////////////////////////////////////////////////////////////////
 // get OpenMP number of procs
 
 // Get the number of processors in the device.
 INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
   if (!isSPMDExecutionMode)
     return GetNumberOfWorkersInTeam();
   return GetNumberOfThreadsInBlock();
 }

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), /*isRuntimeUninitialized=*/false);
+    int numberOfActiveOMPThreads =
+        GetNumberOfOmpThreads(tid, 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);
 

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/nested.c

 // RUN: %compile-run-and-check
 
 #include <omp.h>
 #include <stdio.h>
 
 const int MaxThreads = 1024;
 const int NumThreads = 64;
+const int NumThreads1 = 1;
 
 int main(int argc, char *argv[]) {
   int inParallel = -1, numThreads = -1, threadNum = -1;
   int check1[MaxThreads];
   int check2[MaxThreads];
   for (int i = 0; i < MaxThreads; i++) {
     check1[i] = check2[i] = 0;
   }
 
-  #pragma omp target map(inParallel, numThreads, threadNum, check1[:], check2[:])
   {
     inParallel = omp_in_parallel();
     numThreads = omp_get_num_threads();
     threadNum = omp_get_thread_num();
 
-    // Expecting active parallel region.
-    #pragma omp parallel num_threads(NumThreads)
     {
       int id = omp_get_thread_num();
       check1[id] += omp_get_num_threads() + omp_in_parallel();
 
-      // Expecting serialized parallel region.
-      #pragma omp parallel
       {
         // Expected to be 1.
         int nestedInParallel = omp_in_parallel();
         // Expected to be 1.
         int nestedNumThreads = omp_get_num_threads();
         // Expected to be 0.
         int nestedThreadNum = omp_get_thread_num();
-        #pragma omp atomic
         check2[id] += nestedInParallel + nestedNumThreads + nestedThreadNum;
       }
     }
   }
 
   // CHECK: target: inParallel = 0, numThreads = 1, threadNum = 0
   printf("target: inParallel = %d, numThreads = %d, threadNum = %d\n",
          inParallel, numThreads, threadNum);
 
   // CHECK-NOT: invalid
   for (int i = 0; i < MaxThreads; i++) {
     // Check that all threads reported
     // omp_get_num_threads() = 64, omp_in_parallel() = 1.
     int Expected = NumThreads + 1;
     if (i < NumThreads) {
       if (check1[i] != Expected) {
-        printf("invalid: check1[%d] should be %d, is %d\n", i, Expected, check1[i]);
+        printf("invalid: check1[%d] should be %d, is %d\n", i, Expected,
+               check1[i]);
       }
     } else if (check1[i] != 0) {
       printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
     }
 
     // Check serialized parallel region.
     if (i < NumThreads) {
       if (check2[i] != 2) {
         printf("invalid: check2[%d] should be 2, is %d\n", i, check2[i]);
       }
     } else if (check2[i] != 0) {
       printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
     }
   }
 
+  inParallel = -1;
+  numThreads = -1;
+  threadNum = -1;
+  for (int i = 0; i < MaxThreads; i++) {
+    check1[i] = check2[i] = 0;
+  }
+
+#pragma omp target map(inParallel, numThreads, threadNum, check1[:], check2[:])
+  {
+    inParallel = omp_in_parallel();
+    numThreads = omp_get_num_threads();
+    threadNum = omp_get_thread_num();
+
+// Expecting active parallel region.
+#pragma omp parallel num_threads(NumThreads1)
+    {
+      int id = omp_get_thread_num();
+      check1[id] += omp_get_num_threads() + omp_in_parallel();
+
+// Expecting serialized parallel region.
+#pragma omp parallel
+      {
+        // Expected to be 0.
+        int nestedInParallel = omp_in_parallel();
+        // Expected to be 1.
+        int nestedNumThreads = omp_get_num_threads();
+        // Expected to be 0.
+        int nestedThreadNum = omp_get_thread_num();
+#pragma omp atomic
+        check2[id] += nestedInParallel + nestedNumThreads + nestedThreadNum;
+      }
+    }
+  }
+
+  // CHECK: target: inParallel = 0, numThreads = 1, threadNum = 0
+  printf("target: inParallel = %d, numThreads = %d, threadNum = %d\n",
+         inParallel, numThreads, threadNum);
+
+  // CHECK-NOT: invalid
+  for (int i = 0; i < MaxThreads; i++) {
+    // Check that all threads reported
+    // omp_get_num_threads() = 1, omp_in_parallel() = 0.
+    int Expected = 1;
+    if (i < NumThreads1) {
+      if (check1[i] != Expected) {
+        printf("invalid: check1[%d] should be %d, is %d\n", i, Expected,
+               check1[i]);
+      }
+    } else if (check1[i] != 0) {
+      printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
+    }
+
+    // Check serialized parallel region.
+    if (i < NumThreads1) {
+      if (check2[i] != 1) {
+        printf("invalid: check2[%d] should be 1, is %d\n", i, check2[i]);
+      }
+    } else if (check2[i] != 0) {
+      printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
+    }
+  }
+
   return 0;
 }