[OPENMP][NVPTX] Lightweight runtime support for SPMD mode.

[ABataev]

Summary

Implemented simple and lightweight runtime support for SPMD mode-based
constructs. It adds support for L2 sequential parallelism wihtout full
runtime support. Also, patch fixes some use cases for
uninitialized|lightweight runtime.

[gtbercea]

LGTM

[Hahnfeld]

Great that we don't need any additional entry points. I suppose LLVM is just optimizing away unneeded code when the runtime is inlined?

[ABataev]

In D51222#1212676, @Hahnfeld wrote:

Great that we don't need any additional entry points. I suppose LLVM is just optimizing away unneeded code when the runtime is inlined?

I hope so. We can check later if we can optimize it further.

[Hahnfeld]

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

192–200

Again looking at this code: Shouldn't all threads return, ie should it be outside of if (threadId == 0)?

[ABataev]

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

192–200

You're right, that was an original intention. Could commit a quick fix?

[Hahnfeld]

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

192–200

Yes please.

[ABataev]

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

192–200

I meant could you commit the fix? :) I don't have an access to my laptop.

[Hahnfeld]

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

192–200

Ah, done in r341328. Thanks for your quick replies!

[Hahnfeld]

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

146–167

Has this change related to kmp_sched_static_balanced_chunk slipped into this patch on commit? I can't find it in the original review.

(No need to revert, just curious...)

[ABataev]

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

146–167

Yes, during testing I found some inconsistency between clang and NVPTX libomp abd fixed all of them in a single patch.

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

 
 // Initialize the shared data structures. This is expected to be called for the
 // master thread and warp masters. \param RootS: A pointer to the root of the
 // data sharing stack. \param InitialDataSize: The initial size of the data in
 // the slot.
 EXTERN void
 __kmpc_initialize_data_sharing_environment(__kmpc_data_sharing_slot *rootS,
                                            size_t InitialDataSize) {
-
+  assert(isRuntimeInitialized() && "Expected initialized runtime.");
   DSPRINT0(DSFLAG_INIT,
            "Entering __kmpc_initialize_data_sharing_environment\n");
 
   unsigned WID = getWarpId();
   DSPRINT(DSFLAG_INIT, "Warp ID: %d\n", WID);
 
   omptarget_nvptx_TeamDescr *teamDescr =
       &omptarget_nvptx_threadPrivateContext->TeamContext();
__kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
   return P;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Runtime functions for trunk data sharing scheme.
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE void data_sharing_init_stack_common() {
+  assert(isRuntimeInitialized() && "Expected initialized runtime.");
   omptarget_nvptx_TeamDescr *teamDescr =
       &omptarget_nvptx_threadPrivateContext->TeamContext();
 
   for (int WID = 0; WID < WARPSIZE; WID++) {
     __kmpc_data_sharing_slot *RootS = teamDescr->GetPreallocatedSlotAddr(WID);
     DataSharingState.SlotPtr[WID] = RootS;
     DataSharingState.StackPtr[WID] = (void *)&RootS->Data[0];
   }
 }
 
 // Initialize data sharing data structure. This function needs to be called
 // once at the beginning of a data sharing context (coincides with the kernel
 // initialization). This function is called only by the MASTER thread of each
 // team in non-SPMD mode.
 EXTERN void __kmpc_data_sharing_init_stack() {
+  assert(isRuntimeInitialized() && "Expected initialized runtime.");
   // This function initializes the stack pointer with the pointer to the
   // statically allocated shared memory slots. The size of a shared memory
   // slot is pre-determined to be 256 bytes.
   data_sharing_init_stack_common();
   omptarget_nvptx_globalArgs.Init();
 }
 
 // Initialize data sharing data structure. This function needs to be called
 // once at the beginning of a data sharing context (coincides with the kernel
 // initialization). This function is called in SPMD mode only.
 EXTERN void __kmpc_data_sharing_init_stack_spmd() {
+  assert(isRuntimeInitialized() && "Expected initialized runtime.");
   // This function initializes the stack pointer with the pointer to the
   // statically allocated shared memory slots. The size of a shared memory
   // slot is pre-determined to be 256 bytes.
   if (threadIdx.x == 0)
     data_sharing_init_stack_common();
 
   __threadfence_block();
 }

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

 typedef enum kmp_sched_t {
   kmp_sched_static_chunk = 33,
   kmp_sched_static_nochunk = 34,
   kmp_sched_dynamic = 35,
   kmp_sched_guided = 36,
   kmp_sched_runtime = 37,
   kmp_sched_auto = 38,
 
+  kmp_sched_static_balanced_chunk = 45,
+
   kmp_sched_static_ordered = 65,
   kmp_sched_static_nochunk_ordered = 66,
   kmp_sched_dynamic_ordered = 67,
   kmp_sched_guided_ordered = 68,
   kmp_sched_runtime_ordered = 69,
   kmp_sched_auto_ordered = 70,
 
   kmp_sched_distr_static_chunk = 91,

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

EXTERN double omp_get_wtime(void) {
   unsigned long long nsecs;
   asm("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
   double rc = (double)nsecs * TIMER_PRECISION;
   PRINT(LD_IO, "call omp_get_wtime() returns %g\n", rc);
   return rc;
 }
 
 EXTERN void omp_set_num_threads(int num) {
+  // Ignore it for SPMD mode.
+  if (isSPMDMode())
+    return;
+  assert(isRuntimeInitialized() && "Expected initialized runtime.");
   PRINT(LD_IO, "call omp_set_num_threads(num %d)\n", num);
   if (num <= 0) {
     WARNING0(LW_INPUT, "expected positive num; ignore\n");
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
     currTaskDescr->NThreads() = num;
   }
 }
 
 EXTERN int omp_get_num_threads(void) {
   int tid = GetLogicalThreadIdInBlock();
   int rc = GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized());
   PRINT(LD_IO, "call omp_get_num_threads() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_max_threads(void) {
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    // We're already in parallel region.
+    return 1;  // default is 1 thread avail
+  }
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   int rc = 1; // default is 1 thread avail
   if (!currTaskDescr->InParallelRegion()) {
     // not currently in a parallel region... all are available
     rc = GetNumberOfProcsInTeam();
     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 (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return 0;  // default is 0
+  }
   // per contention group.. meaning threads in current team
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   int rc = currTaskDescr->ThreadLimit();
   PRINT(LD_IO, "call omp_get_thread_limit() return %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_thread_num() {
   int tid = GetLogicalThreadIdInBlock();
   int rc = GetOmpThreadId(tid, isSPMDMode(), isRuntimeUninitialized());
   PRINT(LD_IO, "call omp_get_thread_num() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_get_num_procs(void) {
   int rc = GetNumberOfProcsInDevice();
   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()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    rc = 1;  // SPMD mode is always in parallel.
+  } else {
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
     if (currTaskDescr->InParallelRegion()) {
       rc = 1;
     }
+  }
   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
   // actively decide to put a non-final task into a final one.
   int rc = 1;
   PRINT(LD_IO, "call omp_in_final() returns %d\n", rc);
   return rc;
 }
 
 EXTERN void omp_set_dynamic(int flag) {
   PRINT(LD_IO, "call omp_set_dynamic(%d)\n", flag);
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return;
+  }
 
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   if (flag) {
     currTaskDescr->SetDynamic();
   } else {
     currTaskDescr->ClearDynamic();
   }
 }
 
 EXTERN int omp_get_dynamic(void) {
   int rc = 0;
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return rc;
+  }
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   if (currTaskDescr->IsDynamic()) {
     rc = 1;
   }
   PRINT(LD_IO, "call omp_get_dynamic() returns %d\n", rc);
   return rc;
 }
 

 
 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()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
+  }
   int level = 0;
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   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);
   PRINT(LD_IO, "call omp_get_level() returns %d\n", level);
   return level;
 }
 
 EXTERN int omp_get_active_level(void) {
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return 1;
+  }
   int level = 0; // no active level parallelism
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
   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);
   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()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return level == 1 ? GetThreadIdInBlock() : 0;
+  }
   int rc = 0; // default at level 0
   if (level >= 0) {
     int totLevel = omp_get_level();
     if (level <= totLevel) {
       omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
       int steps = totLevel - level;
       PRINT(LD_IO, "backtrack %d steps\n", steps);
       ASSERT0(LT_FUSSY, currTaskDescr,
if (level >= 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 (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return level == 1 ? GetNumberOfThreadsInBlock() : 1;
+  }
   int rc = 1; // default at level 0
   if (level >= 0) {
     int totLevel = omp_get_level();
     if (level <= totLevel) {
       omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
       int steps = totLevel - level;
       ASSERT0(LT_FUSSY, currTaskDescr,
               "do not expect fct to be called in a non-active thread");
if (level <= totLevel) {
       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()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    *kind = omp_sched_static;
+    *modifier = 1;
+  } else {
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
     *kind = currTaskDescr->GetRuntimeSched();
     *modifier = currTaskDescr->RuntimeChunkSize();
+  }
   PRINT(LD_IO, "call omp_get_schedule returns sched %d and modif %d\n",
         (int)*kind, *modifier);
 }
 
 EXTERN void omp_set_schedule(omp_sched_t kind, int modifier) {
   PRINT(LD_IO, "call omp_set_schedule(sched %d, modif %d)\n", (int)kind,
         modifier);
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() &&
+           "expected SPMD mode only with uninitialized runtime.");
+    return;
+  }
   if (kind >= omp_sched_static && kind < omp_sched_auto) {
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
     currTaskDescr->SetRuntimeSched(kind);
     currTaskDescr->RuntimeChunkSize() = modifier;
     PRINT(LD_IOD, "omp_set_schedule did set sched %d & modif %" PRIu64 "\n",
           (int)currTaskDescr->GetRuntimeSched(),
           currTaskDescr->RuntimeChunkSize());
   }

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

INLINE static void for_static_init(int32_t schedtype, int32_t *plastiter,
 
     // copy
     int lastiter = 0;
     T lb = *plower;
     T ub = *pupper;
     ST stride = *pstride;
     T entityId, numberOfEntities;
     // init
-    switch (schedtype) {
+    switch (SCHEDULE_WITHOUT_MODIFIERS(schedtype)) {
     case kmp_sched_static_chunk: {
       if (chunk > 0) {
         entityId =
             GetOmpThreadId(tid, IsSPMDExecutionMode, IsRuntimeUninitialized);
         numberOfEntities = GetNumberOfOmpThreads(tid, IsSPMDExecutionMode,
                                                  IsRuntimeUninitialized);
         ForStaticChunk(lastiter, lb, ub, stride, chunk, entityId,
                        numberOfEntities);
         break;
       }
     } // note: if chunk <=0, use nochunk
+    case kmp_sched_static_balanced_chunk: {
+      if (chunk > 0) {
+        entityId =
+            GetOmpThreadId(tid, IsSPMDExecutionMode, IsRuntimeUninitialized);
+        numberOfEntities = GetNumberOfOmpThreads(tid, IsSPMDExecutionMode,
+                                                 IsRuntimeUninitialized);
+
+        // round up to make sure the chunk is enough to cover all iterations
+        T tripCount = ub - lb + 1; // +1 because ub is inclusive
+        T span = (tripCount + numberOfEntities - 1) / numberOfEntities;
+        // perform chunk adjustment
+        chunk = (span + chunk - 1) & ~(chunk - 1);
+
+        assert(ub >= lb && "ub must be >= lb.");
+        T oldUb = ub;
+        ForStaticChunk(lastiter, lb, ub, stride, chunk, entityId,
+                       numberOfEntities);
+        if (ub > oldUb)
+          ub = oldUb;
+        break;
+      }
+    } // note: if chunk <=0, use nochunk
     case kmp_sched_static_nochunk: {
       entityId =
           GetOmpThreadId(tid, IsSPMDExecutionMode, IsRuntimeUninitialized);
       numberOfEntities = GetNumberOfOmpThreads(tid, IsSPMDExecutionMode,
                                                IsRuntimeUninitialized);
       ForStaticNoChunk(lastiter, lb, ub, stride, chunk, entityId,
                        numberOfEntities);
       break;
default: {
                      numberOfEntities);
     }
     }
     // copy back
     *plastiter = lastiter;
     *plower = lb;
     *pupper = ub;
     *pstride = stride;
-    PRINT(LD_LOOP,
-          "Got sched: Active %d, total %d: lb %lld, ub %lld, stride %lld\n",
-          GetNumberOfOmpThreads(tid, IsSPMDExecutionMode,
-                                IsRuntimeUninitialized),
-          GetNumberOfWorkersInTeam(), P64(*plower), P64(*pupper),
-          P64(*pstride));
+    PRINT(
+        LD_LOOP,
+        "Got sched: Active %d, total %d: lb %lld, ub %lld, stride %lld, last "
+        "%d\n",
+        GetNumberOfOmpThreads(tid, IsSPMDExecutionMode, IsRuntimeUninitialized),
+        GetNumberOfWorkersInTeam(), P64(*plower), P64(*pupper), P64(*pstride),
+        lastiter);
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for dispatch Init
 
   INLINE static int OrderedSchedule(kmp_sched_t schedule) {
     return schedule >= kmp_sched_ordered_first &&
            schedule <= kmp_sched_ordered_last;
   }
 
   INLINE static void dispatch_init(kmp_Indent *loc, int32_t threadId,
                                    kmp_sched_t schedule, T lb, T ub, ST st,
                                    ST chunk) {
+    assert(isRuntimeInitialized() &&
+           "Expected non-SPMD mode + initialized runtime.");
     int tid = GetLogicalThreadIdInBlock();
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(tid);
     T tnum = currTaskDescr->ThreadsInTeam();
     T tripCount = ub - lb + 1; // +1 because ub is inclusive
     ASSERT0(
         LT_FUSSY,
         GetOmpThreadId(tid, isSPMDMode(), isRuntimeUninitialized()) <
             GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
if (schedule == kmp_sched_static_chunk) {
       omptarget_nvptx_threadPrivateContext->Stride(tid) = stride;
       PRINT(LD_LOOP,
             "dispatch init (static chunk) : num threads = %d, ub =  %" PRId64
             ", next lower bound = %llu, stride = %llu\n",
             GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
             omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid),
             omptarget_nvptx_threadPrivateContext->NextLowerBound(tid),
             omptarget_nvptx_threadPrivateContext->Stride(tid));
+    } else if (schedule == kmp_sched_static_balanced_chunk) {
+      ASSERT0(LT_FUSSY, chunk > 0, "bad chunk value");
+      // save sched state
+      omptarget_nvptx_threadPrivateContext->ScheduleType(tid) = schedule;
+      // save ub
+      omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid) = ub;
+      // compute static chunk
+      ST stride;
+      int lastiter = 0;
+      // round up to make sure the chunk is enough to cover all iterations
+      T span = (tripCount + tnum - 1) / tnum;
+      // perform chunk adjustment
+      chunk = (span + chunk - 1) & ~(chunk - 1);
 
+      T oldUb = ub;
+      ForStaticChunk(
+          lastiter, lb, ub, stride, chunk,
+          GetOmpThreadId(tid, isSPMDMode(), isRuntimeUninitialized()), tnum);
+      assert(ub >= lb && "ub must be >= lb.");
+      if (ub > oldUb)
+        ub = oldUb;
+      // save computed params
+      omptarget_nvptx_threadPrivateContext->Chunk(tid) = chunk;
+      omptarget_nvptx_threadPrivateContext->NextLowerBound(tid) = lb;
+      omptarget_nvptx_threadPrivateContext->Stride(tid) = stride;
+      PRINT(LD_LOOP,
+            "dispatch init (static chunk) : num threads = %d, ub =  %" PRId64
+            ", next lower bound = %llu, stride = %llu\n",
+            GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
+            omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid),
+            omptarget_nvptx_threadPrivateContext->NextLowerBound(tid),
+            omptarget_nvptx_threadPrivateContext->Stride(tid));
     } else if (schedule == kmp_sched_static_nochunk) {
       ASSERT0(LT_FUSSY, chunk == 0, "bad chunk value");
       // save sched state
       omptarget_nvptx_threadPrivateContext->ScheduleType(tid) = schedule;
       // save ub
       omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid) = ub;
       // compute static chunk
       ST stride;
INLINE static int DynamicNextChunk(T &lb, T &ub, T chunkSize,
     return FINISHED;
   }
 
   // On Pascal, with inlining of the runtime into the user application,
   // this code deadlocks.  This is probably because different threads
   // in a warp cannot make independent progress.
   NOINLINE static int dispatch_next(int32_t *plast, T *plower, T *pupper,
                                     ST *pstride) {
+    assert(isRuntimeInitialized() &&
+           "Expected non-SPMD mode + initialized runtime.");
     // ID of a thread in its own warp
 
     // automatically selects thread or warp ID based on selected implementation
     int tid = GetLogicalThreadIdInBlock();
     ASSERT0(
         LT_FUSSY,
         GetOmpThreadId(tid, isSPMDMode(), isRuntimeUninitialized()) <
             GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
NOINLINE static int dispatch_next(int32_t *plast, T *plower, T *pupper,
 
     // not finished (either not finished or last chunk)
     *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\n",
+          "Got sched: active %d, total %d: lb %lld, ub %lld, stride = %lld, "
+          "last %d\n",
           GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()),
-          GetNumberOfWorkersInTeam(), P64(*plower), P64(*pupper),
-          P64(*pstride));
+          GetNumberOfWorkersInTeam(), P64(*plower), P64(*pupper), P64(*pstride),
+          *plast);
     return DISPATCH_NOTFINISHED;
   }
 
   INLINE static void dispatch_fini() {
     // nothing
   }
 
   ////////////////////////////////////////////////////////////////////////////////
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
   }
 #endif
 }
 }; // namespace
 
 EXTERN void __kmpc_reduce_conditional_lastprivate(kmp_Indent *loc, int32_t gtid,
                                                   int32_t varNum, void *array) {
   PRINT0(LD_IO, "call to __kmpc_reduce_conditional_lastprivate(...)\n");
+  assert(isRuntimeInitialized() &&
+         "Expected non-SPMD mode + initialized runtime.");
 
   omptarget_nvptx_TeamDescr &teamDescr = getMyTeamDescriptor();
   int tid = GetOmpThreadId(GetLogicalThreadIdInBlock(), isSPMDMode(),
                            isRuntimeUninitialized());
   uint32_t NumThreads = GetNumberOfOmpThreads(
       GetLogicalThreadIdInBlock(), isSPMDMode(), isRuntimeUninitialized());
   uint64_t *Buffer = teamDescr.getLastprivateIterBuffer();
   for (unsigned i = 0; i < varNum; i++) {

openmp/trunk/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];
+
 // 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;
 
 ////////////////////////////////////////////////////////////////////////////////

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

private:
   uint64_t SourceQueue;
 };
 
 /// Device envrionment data
 struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
+class omptarget_nvptx_SimpleThreadPrivateContext {
+  uint16_t par_level[MAX_THREADS_PER_TEAM];
+public:
+  INLINE void Init() {
+    assert(isSPMDMode() && isRuntimeUninitialized() &&
+           "Expected SPMD + uninitialized runtime modes.");
+    par_level[GetThreadIdInBlock()] = 0;
+  }
+  INLINE void IncParLevel() {
+    assert(isSPMDMode() && isRuntimeUninitialized() &&
+           "Expected SPMD + uninitialized runtime modes.");
+    ++par_level[GetThreadIdInBlock()];
+  }
+  INLINE void DecParLevel() {
+    assert(isSPMDMode() && isRuntimeUninitialized() &&
+           "Expected SPMD + uninitialized runtime modes.");
+    assert(par_level[GetThreadIdInBlock()] > 0 &&
+           "Expected parallel level >0.");
+    --par_level[GetThreadIdInBlock()];
+  }
+  INLINE bool InL2OrHigherParallelRegion() const {
+    assert(isSPMDMode() && isRuntimeUninitialized() &&
+           "Expected SPMD + uninitialized runtime modes.");
+    return par_level[GetThreadIdInBlock()] > 0;
+  }
+  INLINE uint16_t GetParallelLevel() const {
+    assert(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__ __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.
 ////////////////////////////////////////////////////////////////////////////////
 

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

 
 extern __device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
 extern __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
+extern __device__ omptarget_nvptx_Queue<
+    omptarget_nvptx_SimpleThreadPrivateContext, OMP_STATE_COUNT>
+    omptarget_nvptx_device_simpleState[MAX_SM];
+
+extern __device__ __shared__ omptarget_nvptx_SimpleThreadPrivateContext
+    *omptarget_nvptx_simpleThreadPrivateContext;
+
 //
 // The team master sets the outlined function and its arguments in these
 // variables to communicate with the workers.  Since they are in shared memory,
 // there is one copy of these variables for each kernel, instance, and team.
 //
 extern volatile __device__ __shared__ omptarget_nvptx_WorkFn
     omptarget_nvptx_workFn;
 extern __device__ __shared__ uint32_t execution_param;
PRINT(LD_IO, "call to __kmpc_kernel_init_params with version %f\n",
         OMPTARGET_NVPTX_VERSION);
 
   SetTeamsReductionScratchpadPtr(Ptr);
 }
 
 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);
-
-  if (!RequiresOMPRuntime) {
-    // If OMP runtime is not required don't initialize OMP state.
-    setExecutionParameters(Generic, RuntimeUninitialized);
-    return;
-  }
+  assert(RequiresOMPRuntime && "Generic always requires initialized runtime.");
   setExecutionParameters(Generic, RuntimeInitialized);
 
   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.
#endif
   // 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;
 }
 
 EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized) {
-  if (IsOMPRuntimeInitialized) {
-    // Enqueue omp state object for use by another team.
+  assert(IsOMPRuntimeInitialized &&
+         "Generic always requires initialized runtime.");
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
-    int slot = omptarget_nvptx_threadPrivateContext->GetSourceQueue();
 #else
-    int slot = smid() % MAX_SM;
 #endif
-    omptarget_nvptx_device_State[slot].Enqueue(
-        omptarget_nvptx_threadPrivateContext);
-  }
   // 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) {
+      int slot = smid() % MAX_SM;
+      omptarget_nvptx_simpleThreadPrivateContext =
+          omptarget_nvptx_device_simpleState[slot].Dequeue();
+    }
+    __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
EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
   }
 }
 
 EXTERN void __kmpc_spmd_kernel_deinit() {
   // We're not going to pop the task descr stack of each thread since
   // there are no more parallel regions in SPMD mode.
   __syncthreads();
   int threadId = GetThreadIdInBlock();
+  if (isRuntimeUninitialized()) {
+    if (threadId == 0) {
+      // Enqueue omp state object for use by another team.
+      int slot = smid() % MAX_SM;
+      omptarget_nvptx_device_simpleState[slot].Enqueue(
+          omptarget_nvptx_simpleThreadPrivateContext);
+      return;
+    }
+  }
   if (threadId == 0) {
     // Enqueue omp state object for use by another team.
     int slot = smid() % MAX_SM;
     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() {
   return isSPMDMode();
 }

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

 //        blockDim.x;
 //      }
 //    }
 //
 // This routine is always called by the team master..
 EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
                                            int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO, "call to __kmpc_kernel_prepare_parallel\n");
-  omptarget_nvptx_workFn = WorkFn;
+  assert(IsOMPRuntimeInitialized && "expected initialized runtime.");
 
-  if (!IsOMPRuntimeInitialized)
-    return;
+  omptarget_nvptx_workFn = WorkFn;
 
   // This routine is only called by the team master.  The team master is
   // the first thread of the last warp.  It always has the logical thread
   // id of 0 (since it is a shadow for the first worker thread).
   int threadId = 0;
   omptarget_nvptx_TaskDescr *currTaskDescr =
       omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
   ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");

 // 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,
                                    int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_parallel\n");
 
+  assert(IsOMPRuntimeInitialized && "expected initialized runtime.");
+
   // Work function and arguments for L1 parallel region.
   *WorkFn = omptarget_nvptx_workFn;
 
-  if (!IsOMPRuntimeInitialized)
-    return true;
-
   // If this is the termination signal from the master, quit early.
   if (!*WorkFn)
     return false;
 
   // 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();
EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
   }
 
   return isActive;
 }
 
 EXTERN void __kmpc_kernel_end_parallel() {
   // pop stack
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_end_parallel\n");
+  assert(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());
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // support for parallel that goes sequential
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_serialized_parallel(kmp_Indent *loc, uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_serialized_parallel\n");
 
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
+    omptarget_nvptx_simpleThreadPrivateContext->IncParLevel();
+    return;
+  }
+
   // assume this is only called for nested parallel
   int threadId = GetLogicalThreadIdInBlock();
 
   // unlike actual parallel, threads in the same team do not share
   // the workTaskDescr in this case and num threads is fixed to 1
 
   // get current task
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
   currTaskDescr->SaveLoopData();
 
   // allocate new task descriptor and copy value from current one, set prev to
   // it
   omptarget_nvptx_TaskDescr *newTaskDescr =
       (omptarget_nvptx_TaskDescr *)SafeMalloc(sizeof(omptarget_nvptx_TaskDescr),
-                                              (char *)"new seq parallel task");
+                                              "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_Indent *loc,
                                            uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_end_serialized_parallel\n");
 
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
+    omptarget_nvptx_simpleThreadPrivateContext->DecParLevel();
+    return;
+  }
+
   // pop stack
   int threadId = GetLogicalThreadIdInBlock();
   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_Indent *loc, uint32_t global_tid) {
   PRINT0(LD_IO, "call to __kmpc_parallel_level\n");
 
+  if (isRuntimeUninitialized()) {
+    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
+    return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
+  }
+
   int threadId = GetLogicalThreadIdInBlock();
   omptarget_nvptx_TaskDescr *currTaskDescr =
       omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
   if (currTaskDescr->InL2OrHigherParallelRegion())
     return 2;
   else if (currTaskDescr->InParallelRegion())
     return 1;
   else

 
 ////////////////////////////////////////////////////////////////////////////////
 // push params
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_push_num_threads(kmp_Indent *loc, int32_t tid,
                                     int32_t num_threads) {
   PRINT(LD_IO, "call kmpc_push_num_threads %d\n", num_threads);
+  assert(isRuntimeInitialized() && "Runtime must be initialized.");
   tid = GetLogicalThreadIdInBlock();
   omptarget_nvptx_threadPrivateContext->NumThreadsForNextParallel(tid) =
       num_threads;
 }
 
 EXTERN void __kmpc_push_simd_limit(kmp_Indent *loc, int32_t tid,
                                    int32_t simd_limit) {
   PRINT(LD_IO, "call kmpc_push_simd_limit %d\n", simd_limit);
+  assert(isRuntimeInitialized() && "Runtime must be initialized.");
   tid = GetLogicalThreadIdInBlock();
   omptarget_nvptx_threadPrivateContext->SimdLimitForNextSimd(tid) = simd_limit;
 }
 
 // Do nothing. The host guarantees we started the requested number of
 // teams and we only need inspection of gridDim.
 
 EXTERN void __kmpc_push_num_teams(kmp_Indent *loc, int32_t tid,

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

 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode,
                           bool isRuntimeUninitialized) {
   // omp_thread_num
   int rc;
 
   if (isRuntimeUninitialized) {
-    rc = GetThreadIdInBlock();
-    if (!isSPMDExecutionMode && rc >= GetMasterThreadID())
+    assert(isSPMDExecutionMode && "Uninitialized runtime with non-SPMD mode.");
+    // For level 2 parallelism all parallel regions are executed sequentially.
+    if (omptarget_nvptx_simpleThreadPrivateContext
+            ->InL2OrHigherParallelRegion())
       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) {
-    rc = isSPMDExecutionMode ? GetNumberOfThreadsInBlock()
-                             : GetNumberOfThreadsInBlock() - WARPSIZE;
+    assert(isSPMDExecutionMode && "Uninitialized runtime with non-SPMD mode.");
+    // For level 2 parallelism all parallel regions are executed sequentially.
+    if (omptarget_nvptx_simpleThreadPrivateContext
+            ->InL2OrHigherParallelRegion())
+      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();
   }
 
   return rc;

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

EXTERN int32_t __kmpc_cancel_barrier(kmp_Indent *loc_ref, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_cancel_barrier\n");
   __kmpc_barrier(loc_ref, tid);
   PRINT0(LD_SYNC, "completed kmpc_cancel_barrier\n");
   return 0;
 }
 
 EXTERN void __kmpc_barrier(kmp_Indent *loc_ref, int32_t tid) {
   if (isRuntimeUninitialized()) {
-    if (isSPMDMode())
-      __kmpc_barrier_simple_spmd(loc_ref, tid);
-    else
-      __kmpc_barrier_simple_generic(loc_ref, tid);
+    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
   } else {
     tid = GetLogicalThreadIdInBlock();
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(tid);
     int numberOfActiveOMPThreads = GetNumberOfOmpThreads(
         tid, isSPMDMode(), /*isRuntimeUninitialized=*/false);
     if (numberOfActiveOMPThreads > 1) {
       if (isSPMDMode()) {

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

 
 EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Indent *loc, uint32_t global_tid,
                                          kmp_TaskDescr *newKmpTaskDescr,
                                          int32_t depNum, void *depList,
                                          int32_t noAliasDepNum,
                                          void *noAliasDepList) {
   PRINT(LD_IO, "call to __kmpc_omp_task_with_deps(task 0x%llx)\n",
         P64(newKmpTaskDescr));
+  assert(isRuntimeInitialized() && "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(
           newKmpTaskDescr, sizeof(omptarget_nvptx_TaskDescr));
   ASSERT0(LT_FUSSY, &newExplicitTaskDescr->kmpTaskDescr == newKmpTaskDescr,
           "bad assumptions");
   omptarget_nvptx_TaskDescr *newTaskDescr = &newExplicitTaskDescr->taskDescr;
   ASSERT0(LT_FUSSY, (uint64_t)newTaskDescr == (uint64_t)newExplicitTaskDescr,
EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Indent *loc, uint32_t global_tid,
   SafeFree(newExplicitTaskDescr, "explicit task descriptor");
   return 0;
 }
 
 EXTERN void __kmpc_omp_task_begin_if0(kmp_Indent *loc, uint32_t global_tid,
                                       kmp_TaskDescr *newKmpTaskDescr) {
   PRINT(LD_IO, "call to __kmpc_omp_task_begin_if0(task 0x%llx)\n",
         P64(newKmpTaskDescr));
+  assert(isRuntimeInitialized() && "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(
           newKmpTaskDescr, sizeof(omptarget_nvptx_TaskDescr));
   ASSERT0(LT_FUSSY, &newExplicitTaskDescr->kmpTaskDescr == newKmpTaskDescr,
           "bad assumptions");
   omptarget_nvptx_TaskDescr *newTaskDescr = &newExplicitTaskDescr->taskDescr;
   ASSERT0(LT_FUSSY, (uint64_t)newTaskDescr == (uint64_t)newExplicitTaskDescr,
           "bad assumptions");
 
   // 2. push new context: update new task descriptor
   int tid = GetLogicalThreadIdInBlock();
   omptarget_nvptx_TaskDescr *parentTaskDescr = getMyTopTaskDescriptor(tid);
   newTaskDescr->CopyForExplicitTask(parentTaskDescr);
   // set new task descriptor as top
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(tid, newTaskDescr);
   // 3... noting to call... is inline
   // 4 & 5 ... done in complete
 }
 
 EXTERN void __kmpc_omp_task_complete_if0(kmp_Indent *loc, uint32_t global_tid,
                                          kmp_TaskDescr *newKmpTaskDescr) {
   PRINT(LD_IO, "call to __kmpc_omp_task_complete_if0(task 0x%llx)\n",
         P64(newKmpTaskDescr));
+  assert(isRuntimeInitialized() && "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(
           newKmpTaskDescr, sizeof(omptarget_nvptx_TaskDescr));
   ASSERT0(LT_FUSSY, &newExplicitTaskDescr->kmpTaskDescr == newKmpTaskDescr,
           "bad assumptions");
   omptarget_nvptx_TaskDescr *newTaskDescr = &newExplicitTaskDescr->taskDescr;
   ASSERT0(LT_FUSSY, (uint64_t)newTaskDescr == (uint64_t)newExplicitTaskDescr,