[Openmp-commits] [openmp] r340956 - [OPENMP][NVPTX] Replace assert() by ASSERT0() macro, NFC.
Alexey Bataev via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 29 12:22:06 PDT 2018
Author: abataev
Date: Wed Aug 29 12:22:06 2018
New Revision: 340956
URL: http://llvm.org/viewvc/llvm-project?rev=340956&view=rev
Log:
[OPENMP][NVPTX] Replace assert() by ASSERT0() macro, NFC.
Required to fix the buildbots.
Modified:
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/task.cu
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu Wed Aug 29 12:22:06 2018
@@ -79,7 +79,7 @@ __device__ static size_t AlignVal(size_t
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_fram
////////////////////////////////////////////////////////////////////////////////
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_comm
// 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_sta
// 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.
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu Wed Aug 29 12:22:06 2018
@@ -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(voi
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(i
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
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) {
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu Wed Aug 29 12:22:06 2018
@@ -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(uin
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(),
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu Wed Aug 29 12:22:06 2018
@@ -60,7 +60,8 @@ EXTERN void __kmpc_kernel_init_params(vo
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 Threa
}
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();
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h Wed Aug 29 12:22:06 2018
@@ -399,30 +399,30 @@ class omptarget_nvptx_SimpleThreadPrivat
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;
}
};
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu Wed Aug 29 12:22:06 2018
@@ -216,7 +216,7 @@ EXTERN void __kmpc_kernel_end_convergent
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
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
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(k
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_parall
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(km
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(
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_
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;
}
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/supporti.h Wed Aug 29 12:22:06 2018
@@ -101,7 +101,8 @@ INLINE int GetOmpThreadId(int threadId,
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 thr
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())
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu Wed Aug 29 12:22:06 2018
@@ -42,7 +42,8 @@ EXTERN int32_t __kmpc_cancel_barrier(kmp
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();
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/task.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/task.cu?rev=340956&r1=340955&r2=340956&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/task.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/task.cu Wed Aug 29 12:22:06 2018
@@ -81,7 +81,7 @@ EXTERN int32_t __kmpc_omp_task_with_deps
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(km
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_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(
More information about the Openmp-commits
mailing list