Index: libomptarget/deviceRTLs/nvptx/src/libcall.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/libcall.cu +++ libomptarget/deviceRTLs/nvptx/src/libcall.cu @@ -37,7 +37,7 @@ PRINT(LD_IO, "call omp_set_num_threads(num %d)\n", num); if (num <= 0) { WARNING0(LW_INPUT, "expected positive num; ignore\n"); - } else if (parallelLevel[GetWarpId()] == 0) { + } else if (getParallelLevel(GetWarpId()) == 0) { nThreads = num; } } @@ -49,12 +49,13 @@ } EXTERN int omp_get_max_threads(void) { - if (parallelLevel[GetWarpId()] > 0) + int parLevel = getParallelLevel(GetWarpId()); + if (parLevel > 0) // We're already in parallel region. return 1; // default is 1 thread avail // Not currently in a parallel region, return what was set. int rc = 1; - if (parallelLevel[GetWarpId()] == 0) + if (parLevel == 0) rc = nThreads; ASSERT0(LT_FUSSY, rc >= 0, "bad number of threads"); PRINT(LD_IO, "call omp_get_max_threads() return %d\n", rc); @@ -84,7 +85,7 @@ } EXTERN int omp_in_parallel(void) { - int rc = parallelLevel[GetWarpId()] > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0; + int rc = getParallelLevel(GetWarpId()) > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0; PRINT(LD_IO, "call omp_in_parallel() returns %d\n", rc); return rc; } @@ -133,13 +134,13 @@ } EXTERN int omp_get_level(void) { - int level = parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1); + int level = getParallelLevel(GetWarpId()) & (OMP_ACTIVE_PARALLEL_LEVEL - 1); PRINT(LD_IO, "call omp_get_level() returns %d\n", level); return level; } EXTERN int omp_get_active_level(void) { - int level = parallelLevel[GetWarpId()] > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0; + int level = getParallelLevel(GetWarpId()) > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0; PRINT(LD_IO, "call omp_get_active_level() returns %d\n", level) return level; } @@ -149,7 +150,7 @@ return level == 1 ? GetThreadIdInBlock() : 0; int rc = -1; // If level is 0 or all parallel regions are not active - return 0. - unsigned parLevel = parallelLevel[GetWarpId()]; + int parLevel = getParallelLevel(GetWarpId()); if (level == 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL) { int totLevel = omp_get_level(); if (level <= totLevel) { @@ -202,7 +203,7 @@ if (isSPMDMode()) return level == 1 ? GetNumberOfThreadsInBlock() : 1; int rc = -1; - unsigned parLevel = parallelLevel[GetWarpId()]; + int parLevel = getParallelLevel(GetWarpId()); // If level is 0 or all parallel regions are not active - return 1. if (level == 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL) { rc = threadsInTeam; Index: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu +++ libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu @@ -44,7 +44,7 @@ "Generic always requires initialized runtime."); setExecutionParameters(Generic, RuntimeInitialized); for (int I = 0; I < MAX_THREADS_PER_TEAM / WARPSIZE; ++I) - parallelLevel[I] = 0; + setParallelLevel(I, 0); int threadIdInBlock = GetThreadIdInBlock(); ASSERT0(LT_FUSSY, threadIdInBlock == GetMasterThreadID(), @@ -98,11 +98,13 @@ int threadId = GetThreadIdInBlock(); if (threadId == 0) { usedSlotIdx = smid() % MAX_SM; - parallelLevel[0] = - 1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0); + setParallelLevel( + 0, + 1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); } else if (GetLaneId() == 0) { - parallelLevel[GetWarpId()] = - 1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0); + setParallelLevel( + GetWarpId(), + 1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); } if (!RequiresOMPRuntime) { // Runtime is not required - exit. Index: libomptarget/deviceRTLs/nvptx/src/parallel.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/parallel.cu +++ libomptarget/deviceRTLs/nvptx/src/parallel.cu @@ -401,7 +401,7 @@ EXTERN uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid) { PRINT0(LD_IO, "call to __kmpc_parallel_level\n"); - return parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1); + return getParallelLevel(GetWarpId()) & (OMP_ACTIVE_PARALLEL_LEVEL - 1); } // This kmpc call returns the thread id across all teams. It's value is Index: libomptarget/deviceRTLs/nvptx/src/support.h =================================================================== --- libomptarget/deviceRTLs/nvptx/src/support.h +++ libomptarget/deviceRTLs/nvptx/src/support.h @@ -31,6 +31,10 @@ INLINE bool isRuntimeUninitialized(); INLINE bool isRuntimeInitialized(); +INLINE int getParallelLevel(int WarpId); +INLINE void setParallelLevel(int WarpId, int Val); +INLINE void changeParallelLevel(int WarpId, int Val, bool IsIncrement); + //////////////////////////////////////////////////////////////////////////////// // get info from machine //////////////////////////////////////////////////////////////////////////////// Index: libomptarget/deviceRTLs/nvptx/src/supporti.h =================================================================== --- libomptarget/deviceRTLs/nvptx/src/supporti.h +++ libomptarget/deviceRTLs/nvptx/src/supporti.h @@ -84,6 +84,62 @@ return !checkRuntimeUninitialized(loc); } +#ifdef __LP64__ +#define PTR_CONSTRAINT "l" +#else // __LP64__ +#define PTR_CONSTRAINT "r" +#endif // __LP64__ + +INLINE int getParallelLevel(int WarpId) { + // Use volatile access in case of CUDA8 to prevent dangerous optimizations + // performed by ptxas. + int32_t parLevel; +#if defined(CUDA_VERSION) && CUDA_VERSION <= 8000 + asm volatile("ld.volatile.u8 %0, [%1];" + : "=r"(parLevel) + : PTR_CONSTRAINT(¶llelLevel[WarpId]) + : "memory"); +#else + parLevel = parallelLevel[WarpId]; +#endif // CUDA_VERSION + return parLevel; +} + +INLINE void setParallelLevel(int WarpId, int Val) { + uint8_t &parLevel = parallelLevel[WarpId]; +#if defined(CUDA_VERSION) && CUDA_VERSION <= 8000 + asm volatile("st.volatile.u8 [%1], %0;" ::"r"(Val), PTR_CONSTRAINT(&parLevel) + : "memory"); +#else // CUDA_VERSION + parLevel = Val; +#endif // CUDA_VERSION +} + +INLINE void changeParallelLevel(int WarpId, int Val, bool IsIncrement) { + uint8_t &parLevel = parallelLevel[WarpId]; +#if defined(CUDA_VERSION) && CUDA_VERSION <= 8000 + int CurVal; + asm volatile("ld.volatile.u8 %0, [%1];" + : "=r"(CurVal) + : PTR_CONSTRAINT(&parLevel) + : "memory"); + if (IsIncrement) + CurVal += Val; + else + CurVal -= Val; + asm volatile("st.volatile.u8 [%1], %0;" ::"r"(CurVal), + PTR_CONSTRAINT(&parLevel) + : "memory"); +#else // CUDA_VERSION + if (IsIncrement) + parLevel += Val; + else + parLevel -= Val; +#endif // CUDA_VERSION +} + +#undef PTR_CONSTRAINT + //////////////////////////////////////////////////////////////////////////////// // support: get info from machine //////////////////////////////////////////////////////////////////////////////// @@ -152,7 +208,8 @@ INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) { // omp_thread_num 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; } else if (isSPMDExecutionMode) { rc = GetThreadIdInBlock(); @@ -168,7 +225,7 @@ INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode) { // omp_num_threads int rc; - int Level = parallelLevel[GetWarpId()]; + int Level = getParallelLevel(GetWarpId()); if (Level != OMP_ACTIVE_PARALLEL_LEVEL + 1) { rc = 1; } else if (isSPMDExecutionMode) { @@ -206,8 +263,9 @@ int leader = __ffs(tnum) - 1; __SHFL_SYNC(tnum, leader, leader); if (GetLaneId() == leader) { - parallelLevel[GetWarpId()] += - (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); + changeParallelLevel(GetWarpId(), + 1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0), + /*IsIncrement=*/true); } __SHFL_SYNC(tnum, leader, leader); } @@ -217,8 +275,9 @@ int leader = __ffs(tnum) - 1; __SHFL_SYNC(tnum, leader, leader); if (GetLaneId() == leader) { - parallelLevel[GetWarpId()] -= - (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0)); + changeParallelLevel(GetWarpId(), + 1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0), + /*IsIncrement=*/false); } __SHFL_SYNC(tnum, leader, leader); } Index: libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp =================================================================== --- libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp +++ libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp @@ -47,5 +47,15 @@ // CHECK: Expected count = 86 printf("Expected count = %d\n", Count); + Count = 0; +#pragma omp target parallel for reduction(+: Count) schedule(dynamic, 2) num_threads(64) + for (int J = 0; J < 1000; ++J) { + Count += J; + } + + // Final result of Count is 1000 * (999-0) / 2 + // CHECK: Expected count with dynamic scheduling = 499500 + printf("Expected count with dynamic scheduling = %d\n", Count); + return isHost; }