From d256797c9035aebf0309489c04dc34f8bae49dc4 Mon Sep 17 00:00:00 2001 From: JonChesterfield Date: Tue, 29 Sep 2020 23:11:46 +0100 Subject: [PATCH] [nfc][libomptarget] Drop parameter to named_sync [nfc][libomptarget] Drop parameter to named_sync named_sync has one call site (in sync.cu) where it always passed L1_BARRIER. Folding this into the call site and dropping the macro is a simplification. amdgpu doesn't have ptx' bar.sync instruction. A correct implementation of __kmpc_impl_named_sync in terms of shared memory is much easier if it can assume that the barrier argument is this constant. Said implementation is left for a second patch. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D88474 --- openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h | 11 +++-------- openmp/libomptarget/deviceRTLs/common/src/sync.cu | 3 +-- openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h | 9 ++++----- 3 files changed, 8 insertions(+), 15 deletions(-) diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h index 3c90b39..8afc5e7 100644 --- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h @@ -42,10 +42,6 @@ #define WARPSIZE 64 -// The named barrier for active parallel threads of a team in an L1 parallel -// region to synchronize with each other. -#define L1_BARRIER (1) - // Maximum number of preallocated arguments to an outlined parallel/simd // function. Anything more requires dynamic memory allocation. #define MAX_SHARED_ARGS 20 @@ -113,10 +109,9 @@ INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) { // AMDGCN doesn't need to sync threads in a warp } -INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) { - // we have protected the master warp from releasing from its barrier - // due to a full workgroup barrier in the middle of a work function. - // So it is ok to issue a full workgroup barrier here. +INLINE void __kmpc_impl_named_sync(uint32_t num_threads) { + (void)num_threads; + // TODO: Implement on top of __SHARED__ __builtin_amdgcn_s_barrier(); } diff --git a/openmp/libomptarget/deviceRTLs/common/src/sync.cu b/openmp/libomptarget/deviceRTLs/common/src/sync.cu index 3979e20..824094c 100644 --- a/openmp/libomptarget/deviceRTLs/common/src/sync.cu +++ b/openmp/libomptarget/deviceRTLs/common/src/sync.cu @@ -60,8 +60,7 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) { PRINT(LD_SYNC, "call kmpc_barrier with %d omp threads, sync parameter %d\n", (int)numberOfActiveOMPThreads, (int)threads); - // Barrier #1 is for synchronization among active threads. - __kmpc_impl_named_sync(L1_BARRIER, threads); + __kmpc_impl_named_sync(threads); } } else { // Still need to flush the memory per the standard. diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h index e3a3d0f..f7bc7e1 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h @@ -37,10 +37,6 @@ #define WARPSIZE 32 -// The named barrier for active parallel threads of a team in an L1 parallel -// region to synchronize with each other. -#define L1_BARRIER (1) - // Maximum number of preallocated arguments to an outlined parallel/simd function. // Anything more requires dynamic memory allocation. #define MAX_SHARED_ARGS 20 @@ -187,7 +183,10 @@ INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) { #endif // CUDA_VERSION } -INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) { +INLINE void __kmpc_impl_named_sync(uint32_t num_threads) { + // The named barrier for active parallel threads of a team in an L1 parallel + // region to synchronize with each other. + int barrier = 1; asm volatile("bar.sync %0, %1;" : : "r"(barrier), "r"(num_threads) -- 2.7.4