[Openmp-commits] [openmp] 56adceb - [libomptarget][nfc] Add nop syncwarp function for amdgcn

[Jon Chesterfield via Openmp-commits]

Author: Jon Chesterfield
Date: 2019-12-13T14:27:52Z
New Revision: 56adcebfda84009b2825e69f71068685360abed7

URL: https://github.com/llvm/llvm-project/commit/56adcebfda84009b2825e69f71068685360abed7
DIFF: https://github.com/llvm/llvm-project/commit/56adcebfda84009b2825e69f71068685360abed7.diff

LOG: [libomptarget][nfc] Add nop syncwarp function for amdgcn

Added: 
    

Modified: 
    openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
index d9b61fd75581..5082d469d050 100644
--- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
@@ -121,6 +121,10 @@ EXTERN int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
 
 INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
+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.