Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-12-21 06:52:10 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[libomptarget] Refactor syncwarp macro to inline function

Browse Source 
Summary:
[libomptarget] Refactor syncwarp macro to inline function
See also abandoned D66846, split into this diff and others.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66857

llvm-svn: 370149
This commit is contained in:
Jon Chesterfield 2019-08-28 02:02:53 +00:00
parent e73e3013a6
commit be3d487313
3 changed files with 13 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
View File

@ -52,11 +52,8 @@
#error CUDA_VERSION macro is undefined, something wrong with cuda.
#elif CUDA_VERSION >= 9000
#define __ACTIVEMASK() __activemask()
#define __SYNCWARP(Mask) __syncwarp(Mask)
#else
#define __ACTIVEMASK() __ballot(1)
// In Cuda < 9.0 no need to sync threads in warps.
#define __SYNCWARP(Mask)
#endif // CUDA_VERSION
#define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");

10
openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
View File

@ -14,6 +14,8 @@
// Execution Parameters
////////////////////////////////////////////////////////////////////////////////
#include "target_impl.h"
INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
execution_param = EMode;
execution_param |= RMode;
@ -203,7 +205,7 @@ INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
INLINE void IncParallelLevel(bool ActiveParallel) {
unsigned Active = __ACTIVEMASK();
__SYNCWARP(Active);
__kmpc_impl_syncwarp(Active);
unsigned LaneMaskLt;
asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
unsigned Rank = __popc(Active & LaneMaskLt);
@ -212,12 +214,12 @@ INLINE void IncParallelLevel(bool ActiveParallel) {
(1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
__threadfence();
}
__SYNCWARP(Active);
__kmpc_impl_syncwarp(Active);
}
INLINE void DecParallelLevel(bool ActiveParallel) {
unsigned Active = __ACTIVEMASK();
__SYNCWARP(Active);
__kmpc_impl_syncwarp(Active);
unsigned LaneMaskLt;
asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
unsigned Rank = __popc(Active & LaneMaskLt);
@ -226,7 +228,7 @@ INLINE void DecParallelLevel(bool ActiveParallel) {
(1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
__threadfence();
}
__SYNCWARP(Active);
__kmpc_impl_syncwarp(Active);
}
////////////////////////////////////////////////////////////////////////////////

8
openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
View File

@ -63,6 +63,12 @@ INLINE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t Mask,
#endif // CUDA_VERSION
}
INLINE void __kmpc_impl_syncwarp(int32_t Mask) { __SYNCWARP(Mask); }
INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
#if CUDA_VERSION >= 9000
__syncwarp(Mask);
#else
// In Cuda < 9.0 no need to sync threads in warps.
#endif // CUDA_VERSION
}
#endif