Changeset View
Changeset View
Standalone View
Standalone View
libomptarget/deviceRTLs/nvptx/src/supporti.h
| Show First 20 Lines • Show All 80 Lines • ▼ Show 20 Lines | INLINE bool checkRuntimeUninitialized(kmp_Ident *loc) { | ||||
| // functions. | // functions. | ||||
| return isRuntimeUninitialized(); | return isRuntimeUninitialized(); | ||||
| } | } | ||||
| INLINE bool checkRuntimeInitialized(kmp_Ident *loc) { | INLINE bool checkRuntimeInitialized(kmp_Ident *loc) { | ||||
| return !checkRuntimeUninitialized(loc); | return !checkRuntimeUninitialized(loc); | ||||
| } | } | ||||
| INLINE int getParallelLevel(int WarpId) { | |||||
| return __kmpc_impl_get_parallel_level(parallelLevel[WarpId]); | |||||
| } | |||||
| INLINE void setParallelLevel(int WarpId, int Val) { | |||||
| __kmpc_impl_set_parallel_level(parallelLevel[WarpId], Val); | |||||
| } | |||||
| INLINE void changeParallelLevel(int WarpId, int Val, bool IsIncrement) { | |||||
| int ParLevel = getParallelLevel(WarpId); | |||||
| if (IsIncrement) | |||||
| ParLevel += Val; | |||||
| else | |||||
tra: Perhaps it can be deduped into something like this:
```
#ifdef __LP64__
#define PTR_CONSTRAINT… | |||||
Sure, thanks. ABataev: Sure, thanks. | |||||
| ParLevel -= Val; | |||||
| setParallelLevel(WarpId, ParLevel); | |||||
| } | |||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| // support: get info from machine | // support: get info from machine | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| // | // | ||||
| // Calls to the NVPTX layer (assuming 1D layout) | // Calls to the NVPTX layer (assuming 1D layout) | ||||
| // | // | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| INLINE int GetThreadIdInBlock() { return threadIdx.x; } | INLINE int GetThreadIdInBlock() { return threadIdx.x; } | ||||
| INLINE int GetBlockIdInKernel() { return blockIdx.x; } | INLINE int GetBlockIdInKernel() { return blockIdx.x; } | ||||
| INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; } | INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; } | ||||
Could this function call getParallelLevel, increment/decrement the result, then call setParallelLevel? JonChesterfield: Could this function call getParallelLevel, increment/decrement the result, then call… | |||||
| INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; } | INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; } | ||||
| INLINE unsigned GetWarpId() { return threadIdx.x / WARPSIZE; } | INLINE unsigned GetWarpId() { return threadIdx.x / WARPSIZE; } | ||||
| INLINE unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); } | INLINE unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); } | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| // | // | ||||
| Show All 36 Lines | |||||
| // | // | ||||
| // OpenMP Thread Support Layer | // OpenMP Thread Support Layer | ||||
| // | // | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) { | INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) { | ||||
| // omp_thread_num | // omp_thread_num | ||||
| int rc; | int rc; | ||||
| if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) { | int Level = getParallelLevel(GetWarpId()); | ||||
| if ((Level & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) { | |||||
| rc = 0; | rc = 0; | ||||
| } else if (isSPMDExecutionMode) { | } else if (isSPMDExecutionMode) { | ||||
| rc = GetThreadIdInBlock(); | rc = GetThreadIdInBlock(); | ||||
| } else { | } else { | ||||
| omptarget_nvptx_TaskDescr *currTaskDescr = | omptarget_nvptx_TaskDescr *currTaskDescr = | ||||
| omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId); | omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId); | ||||
| ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr"); | ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr"); | ||||
| rc = currTaskDescr->ThreadId(); | rc = currTaskDescr->ThreadId(); | ||||
| } | } | ||||
| return rc; | return rc; | ||||
| } | } | ||||
| INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode) { | INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode) { | ||||
| // omp_num_threads | // omp_num_threads | ||||
| int rc; | int rc; | ||||
| int Level = parallelLevel[GetWarpId()]; | int Level = getParallelLevel(GetWarpId()); | ||||
| if (Level != OMP_ACTIVE_PARALLEL_LEVEL + 1) { | if (Level != OMP_ACTIVE_PARALLEL_LEVEL + 1) { | ||||
| rc = 1; | rc = 1; | ||||
| } else if (isSPMDExecutionMode) { | } else if (isSPMDExecutionMode) { | ||||
| rc = GetNumberOfThreadsInBlock(); | rc = GetNumberOfThreadsInBlock(); | ||||
| } else { | } else { | ||||
| rc = threadsInTeam; | rc = threadsInTeam; | ||||
| } | } | ||||
| Show All 21 Lines | |||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| // Parallel level | // Parallel level | ||||
| INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) { | INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) { | ||||
| __kmpc_impl_syncwarp(Mask); | __kmpc_impl_syncwarp(Mask); | ||||
| __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt(); | __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt(); | ||||
| unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt); | unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt); | ||||
| if (Rank == 0) { | if (Rank == 0) { | ||||
| parallelLevel[GetWarpId()] += | changeParallelLevel(GetWarpId(), | ||||
| (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); | 1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0), | ||||
| /*IsIncrement=*/true); | |||||
| __threadfence(); | __threadfence(); | ||||
| } | } | ||||
| __kmpc_impl_syncwarp(Mask); | __kmpc_impl_syncwarp(Mask); | ||||
| } | } | ||||
| INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) { | INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) { | ||||
| __kmpc_impl_syncwarp(Mask); | __kmpc_impl_syncwarp(Mask); | ||||
| __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt(); | __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt(); | ||||
| unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt); | unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt); | ||||
| if (Rank == 0) { | if (Rank == 0) { | ||||
| parallelLevel[GetWarpId()] -= | changeParallelLevel(GetWarpId(), | ||||
| (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); | 1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0), | ||||
| /*IsIncrement=*/false); | |||||
| __threadfence(); | __threadfence(); | ||||
| } | } | ||||
| __kmpc_impl_syncwarp(Mask); | __kmpc_impl_syncwarp(Mask); | ||||
| } | } | ||||
| //////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////// | ||||
| // get OpenMP number of procs | // get OpenMP number of procs | ||||
| ▲ Show 20 Lines • Show All 64 Lines • Show Last 20 Lines | |||||
Perhaps it can be deduped into something like this:
#ifdef __LP64__ #define PTR_CONSTRAINT "l" #else // __LP64__ #define PTR_CONSTRAINT "r" #endif // __LP64__ ... asm volatile("st.volatile.u8 [%1], %0;" ::"r"(Val), : PTR_CONSTRAINT (&parLevel) : "memory");