diff options
| author | Alexey Bataev <a.bataev@hotmail.com> | 2019-05-24 14:06:47 +0000 |
|---|---|---|
| committer | Alexey Bataev <a.bataev@hotmail.com> | 2019-05-24 14:06:47 +0000 |
| commit | 34def82866d4eb9c52a0ae65aa2abefbc4a923c5 (patch) | |
| tree | a349541443c5aaab055383ae1eaf93f8bef567a4 /libomptarget | |
| parent | 926938a039dbca89d50348b06cf7cda4924c60f8 (diff) | |
Revert "[OPENMP][NVPTX]Fix barriers and parallel level counters, NFC."
This reverts commit r361421 to split the patch into 3 parts.
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@361638 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'libomptarget')
| -rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/omp_data.cu | 3 | ||||
| -rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h | 2 | ||||
| -rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/sync.cu | 4 |
3 files changed, 3 insertions, 6 deletions
diff --git a/libomptarget/deviceRTLs/nvptx/src/omp_data.cu b/libomptarget/deviceRTLs/nvptx/src/omp_data.cu index 0cd9b57..d369da1 100644 --- a/libomptarget/deviceRTLs/nvptx/src/omp_data.cu +++ b/libomptarget/deviceRTLs/nvptx/src/omp_data.cu @@ -31,8 +31,7 @@ __device__ omptarget_nvptx_SimpleMemoryManager __device__ __shared__ uint32_t usedMemIdx; __device__ __shared__ uint32_t usedSlotIdx; -__device__ __shared__ volatile uint8_t - parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE]; +__device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE]; __device__ __shared__ uint16_t threadLimit; __device__ __shared__ uint16_t threadsInTeam; __device__ __shared__ uint16_t nThreads; diff --git a/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h index b85d0a7..cd51538 100644 --- a/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h +++ b/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h @@ -398,7 +398,7 @@ extern __device__ omptarget_nvptx_SimpleMemoryManager omptarget_nvptx_simpleMemoryManager; extern __device__ __shared__ uint32_t usedMemIdx; extern __device__ __shared__ uint32_t usedSlotIdx; -extern __device__ __shared__ volatile uint8_t +extern __device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE]; extern __device__ __shared__ uint16_t threadLimit; extern __device__ __shared__ uint16_t threadsInTeam; diff --git a/libomptarget/deviceRTLs/nvptx/src/sync.cu b/libomptarget/deviceRTLs/nvptx/src/sync.cu index 191b046..d81aa8f 100644 --- a/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ b/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -62,8 +62,6 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) { // Barrier #1 is for synchronization among active threads. named_sync(L1_BARRIER, threads); } - } else { - __kmpc_flush(loc_ref); } // numberOfActiveOMPThreads > 1 PRINT0(LD_SYNC, "completed kmpc_barrier\n"); } @@ -132,7 +130,7 @@ EXTERN void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid) { EXTERN void __kmpc_flush(kmp_Ident *loc) { PRINT0(LD_IO, "call kmpc_flush\n"); - __threadfence(); + __threadfence_system(); } //////////////////////////////////////////////////////////////////////////////// |