[OPENMP][NVPTX] Replace assert() by ASSERT0() macro, NFC.

Required to fix the buildbots.

9 files changed, 71 insertions, 64 deletions

diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
index a1b60429935..4a28a7c19d6 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
@@ -79,7 +79,7 @@ __device__ static size_t AlignVal(size_t Val) {
 EXTERN void
 __kmpc_initialize_data_sharing_environment(__kmpc_data_sharing_slot *rootS,
                                            size_t InitialDataSize) {
-  assert(isRuntimeInitialized() && "Expected initialized runtime.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   DSPRINT0(DSFLAG_INIT,
            "Entering __kmpc_initialize_data_sharing_environment\n");
 
@@ -331,7 +331,7 @@ __kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE void data_sharing_init_stack_common() {
-  assert(isRuntimeInitialized() && "Expected initialized runtime.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   omptarget_nvptx_TeamDescr *teamDescr =
       &omptarget_nvptx_threadPrivateContext->TeamContext();
 
@@ -347,7 +347,7 @@ INLINE void data_sharing_init_stack_common() {
 // 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // 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.
@@ -359,7 +359,7 @@ EXTERN void __kmpc_data_sharing_init_stack() {
 // 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // 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.
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
index de493a5f25a..f2599a4e9c2 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
@@ -34,7 +34,7 @@ EXTERN void omp_set_num_threads(int num) {
   // Ignore it for SPMD mode.
   if (isSPMDMode())
     return;
-  assert(isRuntimeInitialized() && "Expected initialized runtime.");
+  ASSERT0(LT_FUSSY, 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");
@@ -53,8 +53,8 @@ EXTERN int omp_get_num_threads(void) {
 
 EXTERN int omp_get_max_threads(void) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     // We're already in parallel region.
     return 1;  // default is 1 thread avail
   }
@@ -71,8 +71,8 @@ EXTERN int omp_get_max_threads(void) {
 
 EXTERN int omp_get_thread_limit(void) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return 0;  // default is 0
   }
   // per contention group.. meaning threads in current team
@@ -98,8 +98,8 @@ EXTERN int omp_get_num_procs(void) {
 EXTERN int omp_in_parallel(void) {
   int rc = 0;
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    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();
@@ -124,8 +124,8 @@ EXTERN int omp_in_final(void) {
 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.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return;
   }
 
@@ -140,8 +140,8 @@ EXTERN void omp_set_dynamic(int flag) {
 EXTERN int omp_get_dynamic(void) {
   int rc = 0;
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return rc;
   }
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
@@ -177,8 +177,8 @@ EXTERN int omp_get_max_active_levels(void) {
 
 EXTERN int omp_get_level(void) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
   }
   int level = 0;
@@ -197,8 +197,8 @@ EXTERN int omp_get_level(void) {
 
 EXTERN int omp_get_active_level(void) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return 1;
   }
   int level = 0; // no active level parallelism
@@ -219,8 +219,8 @@ EXTERN int omp_get_active_level(void) {
 
 EXTERN int omp_get_ancestor_thread_num(int level) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return level == 1 ? GetThreadIdInBlock() : 0;
   }
   int rc = 0; // default at level 0
@@ -267,8 +267,8 @@ EXTERN int omp_get_ancestor_thread_num(int level) {
 
 EXTERN int omp_get_team_size(int level) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return level == 1 ? GetNumberOfThreadsInBlock() : 1;
   }
   int rc = 1; // default at level 0
@@ -299,8 +299,8 @@ EXTERN int omp_get_team_size(int level) {
 
 EXTERN void omp_get_schedule(omp_sched_t *kind, int *modifier) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() &&
-           "expected SPMD mode only with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     *kind = omp_sched_static;
     *modifier = 1;
   } else {
@@ -316,8 +316,8 @@ 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.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode only with uninitialized runtime.");
     return;
   }
   if (kind >= omp_sched_static && kind < omp_sched_auto) {
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu
index c4c8e712adc..bd84f0fad13 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu
@@ -156,7 +156,7 @@ public:
         // perform chunk adjustment
         chunk = (span + chunk - 1) & ~(chunk - 1);
 
-        assert(ub >= lb && "ub must be >= lb.");
+        ASSERT0(LT_FUSSY, ub >= lb, "ub must be >= lb.");
         T oldUb = ub;
         ForStaticChunk(lastiter, lb, ub, stride, chunk, entityId,
                        numberOfEntities);
@@ -241,8 +241,8 @@ public:
   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.");
+    ASSERT0(LT_FUSSY, isRuntimeInitialized(),
+            "Expected non-SPMD mode + initialized runtime.");
     int tid = GetLogicalThreadIdInBlock();
     omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(tid);
     T tnum = currTaskDescr->ThreadsInTeam();
@@ -351,7 +351,7 @@ public:
       ForStaticChunk(
           lastiter, lb, ub, stride, chunk,
           GetOmpThreadId(tid, isSPMDMode(), isRuntimeUninitialized()), tnum);
-      assert(ub >= lb && "ub must be >= lb.");
+      ASSERT0(LT_FUSSY, ub >= lb, "ub must be >= lb.");
       if (ub > oldUb)
         ub = oldUb;
       // save computed params
@@ -454,8 +454,8 @@ public:
   // 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.");
+    ASSERT0(LT_FUSSY, 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
@@ -795,8 +795,8 @@ INLINE void syncWorkersInGenericMode(uint32_t NumThreads) {
 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(),
+          "Expected non-SPMD mode + initialized runtime.");
 
   omptarget_nvptx_TeamDescr &teamDescr = getMyTeamDescriptor();
   int tid = GetOmpThreadId(GetLogicalThreadIdInBlock(), isSPMDMode(),
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
index 884a27109a3..0fe910d23b8 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
@@ -60,7 +60,8 @@ EXTERN void __kmpc_kernel_init_params(void *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);
-  assert(RequiresOMPRuntime && "Generic always requires initialized runtime.");
+  ASSERT0(LT_FUSSY, RequiresOMPRuntime,
+          "Generic always requires initialized runtime.");
   setExecutionParameters(Generic, RuntimeInitialized);
 
   int threadIdInBlock = GetThreadIdInBlock();
@@ -97,8 +98,8 @@ EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime) {
 }
 
 EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized) {
-  assert(IsOMPRuntimeInitialized &&
-         "Generic always requires initialized runtime.");
+  ASSERT0(LT_FUSSY, IsOMPRuntimeInitialized,
+          "Generic always requires initialized runtime.");
   // Enqueue omp state object for use by another team.
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
   int slot = omptarget_nvptx_threadPrivateContext->GetSourceQueue();
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
index 9e9f9345dd3..7058d6050f4 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -399,30 +399,30 @@ class omptarget_nvptx_SimpleThreadPrivateContext {
   uint16_t par_level[MAX_THREADS_PER_TEAM];
 public:
   INLINE void Init() {
-    assert(isSPMDMode() && isRuntimeUninitialized() &&
-           "Expected SPMD + uninitialized runtime modes.");
+    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
+            "Expected SPMD + uninitialized runtime modes.");
     par_level[GetThreadIdInBlock()] = 0;
   }
   INLINE void IncParLevel() {
-    assert(isSPMDMode() && isRuntimeUninitialized() &&
-           "Expected SPMD + uninitialized runtime modes.");
+    ASSERT0(LT_FUSSY, 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.");
+    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 {
-    assert(isSPMDMode() && isRuntimeUninitialized() &&
-           "Expected SPMD + uninitialized runtime modes.");
+    ASSERT0(LT_FUSSY, 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.");
+    ASSERT0(LT_FUSSY, isSPMDMode() && isRuntimeUninitialized(),
+            "Expected SPMD + uninitialized runtime modes.");
     return par_level[GetThreadIdInBlock()] + 1;
   }
 };
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/parallel.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/parallel.cu
index 81d13a4a942..6157aa20794 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/parallel.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/parallel.cu
@@ -216,7 +216,7 @@ EXTERN void __kmpc_kernel_end_convergent_parallel(void *buffer) {
 EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
                                            int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO, "call to __kmpc_kernel_prepare_parallel\n");
-  assert(IsOMPRuntimeInitialized && "expected initialized runtime.");
+  ASSERT0(LT_FUSSY, IsOMPRuntimeInitialized, "Expected initialized runtime.");
 
   omptarget_nvptx_workFn = WorkFn;
 
@@ -319,7 +319,7 @@ 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.");
+  ASSERT0(LT_FUSSY, IsOMPRuntimeInitialized, "Expected initialized runtime.");
 
   // Work function and arguments for L1 parallel region.
   *WorkFn = omptarget_nvptx_workFn;
@@ -361,7 +361,7 @@ EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
 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.");
+  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.
@@ -379,7 +379,8 @@ 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.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode with uninitialized runtime.");
     omptarget_nvptx_simpleThreadPrivateContext->IncParLevel();
     return;
   }
@@ -417,7 +418,8 @@ EXTERN void __kmpc_end_serialized_parallel(kmp_Indent *loc,
   PRINT0(LD_IO, "call to __kmpc_end_serialized_parallel\n");
 
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode with uninitialized runtime.");
     omptarget_nvptx_simpleThreadPrivateContext->DecParLevel();
     return;
   }
@@ -438,7 +440,8 @@ 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.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode with uninitialized runtime.");
     return omptarget_nvptx_simpleThreadPrivateContext->GetParallelLevel();
   }
 
@@ -468,7 +471,7 @@ EXTERN int32_t __kmpc_global_thread_num(kmp_Indent *loc) {
 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   tid = GetLogicalThreadIdInBlock();
   omptarget_nvptx_threadPrivateContext->NumThreadsForNextParallel(tid) =
       num_threads;
@@ -477,7 +480,7 @@ EXTERN void __kmpc_push_num_threads(kmp_Indent *loc, int32_t tid,
 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   tid = GetLogicalThreadIdInBlock();
   omptarget_nvptx_threadPrivateContext->SimdLimitForNextSimd(tid) = simd_limit;
 }
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h b/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
index c58bfc60358..9cdcc162dd4 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
@@ -101,7 +101,8 @@ INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode,
   int rc;
 
   if (isRuntimeUninitialized) {
-    assert(isSPMDExecutionMode && "Uninitialized runtime with non-SPMD mode.");
+    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())
@@ -122,7 +123,8 @@ INLINE int GetNumberOfOmpThreads(int threadId, bool isSPMDExecutionMode,
   int rc;
 
   if (isRuntimeUninitialized) {
-    assert(isSPMDExecutionMode && "Uninitialized runtime with non-SPMD mode.");
+    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())
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu
index ba11f6257b6..0a99405e778 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -42,7 +42,8 @@ EXTERN int32_t __kmpc_cancel_barrier(kmp_Indent *loc_ref, int32_t tid) {
 
 EXTERN void __kmpc_barrier(kmp_Indent *loc_ref, int32_t tid) {
   if (isRuntimeUninitialized()) {
-    assert(isSPMDMode() && "Expected SPMD mode with uninitialized runtime.");
+    ASSERT0(LT_FUSSY, isSPMDMode(),
+            "Expected SPMD mode with uninitialized runtime.");
     __kmpc_barrier_simple_spmd(loc_ref, tid);
   } else {
     tid = GetLogicalThreadIdInBlock();
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/task.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/task.cu
index a8895340c28..f0431abf24d 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/task.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/task.cu
@@ -81,7 +81,7 @@ EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Indent *loc, uint32_t global_tid,
                                          void *noAliasDepList) {
   PRINT(LD_IO, "call to __kmpc_omp_task_with_deps(task 0x%llx)\n",
         P64(newKmpTaskDescr));
-  assert(isRuntimeInitialized() && "Runtime must be initialized.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(
@@ -118,7 +118,7 @@ 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(
@@ -143,7 +143,7 @@ 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.");
+  ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // 1. get explict task descr from kmp task descr
   omptarget_nvptx_ExplicitTaskDescr *newExplicitTaskDescr =
       (omptarget_nvptx_ExplicitTaskDescr *)SUB_BYTES(