diff options
| author | Jon Chesterfield <jonathanchesterfield@gmail.com> | 2019-12-13 14:27:52 +0000 |
|---|---|---|
| committer | Jon Chesterfield <jonathanchesterfield@gmail.com> | 2019-12-13 14:27:52 +0000 |
| commit | 56adcebfda84009b2825e69f71068685360abed7 (patch) | |
| tree | b8f781879ce5a81db685872719806b0d62f6b5bc | |
| parent | dc9e6ba90bebe72f846e76fcc3f2c5145df24613 (diff) | |
| download | bcm5719-llvm-56adcebfda84009b2825e69f71068685360abed7.tar.gz bcm5719-llvm-56adcebfda84009b2825e69f71068685360abed7.zip | |
[libomptarget][nfc] Add nop syncwarp function for amdgcn
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h | 4 |
1 files changed, 4 insertions, 0 deletions
diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h index d9b61fd7558..5082d469d05 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. |
