Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

[OpenMP] Small fixes after refactoring and adding SPIR-V support

Browse source
This commit is contained in:
Joseph Huber 2026-01-27 15:55:36 -06:00
parent 93ed74c890
commit 5b98c9be17
3 changed files with 50 additions and 46 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
openmp/device/CMakeLists.txt
View file

@ -53,8 +53,8 @@ elseif("${LLVM_DEFAULT_TARGET_TRIPLE}" MATCHES "^nvptx" OR
"${CMAKE_CXX_COMPILER_TARGET}" MATCHES "^nvptx")
set(target_name "nvptx")
list(APPEND compile_options --cuda-feature=+ptx63)
elseif("${LLVM_DEFAULT_TARGET_TRIPLE}" MATCHES "^spirv64" OR
"${CMAKE_CXX_COMPILER_TARGET}" MATCHES "^spirv64")
elseif("${LLVM_DEFAULT_TARGET_TRIPLE}" MATCHES "^spirv" OR
"${CMAKE_CXX_COMPILER_TARGET}" MATCHES "^spirv")
set(target_name "spirv")
endif()

84
openmp/device/include/Synchronization.h
View file

@ -192,43 +192,6 @@ atomicExchange(uint32_t *Address, uint32_t Val, atomic::OrderingTy Ordering,
} // namespace atomic
namespace synchronize {
/// Initialize the synchronization machinery. Must be called by all threads.
void init(bool IsSPMD);
/// Synchronize all threads in a warp identified by \p Mask.
static inline void warp(LaneMaskTy Mask) { __gpu_sync_lane(Mask); }
/// Synchronize all threads in a block and perform a fence before and after the
/// barrier according to \p Ordering. Note that the fence might be part of the
/// barrier.
static inline void threads(atomic::OrderingTy Ordering) {
#ifdef __NVPTX__
__nvvm_barrier_sync(8);
#else
__gpu_sync_threads();
#endif
}
/// Synchronizing threads is allowed even if they all hit different instances of
/// `synchronize::threads()`. However, `synchronize::threadsAligned()` is more
/// restrictive in that it requires all threads to hit the same instance. The
/// noinline is removed by the openmp-opt pass and helps to preserve the
/// information till then.
///{
/// Synchronize all threads in a block, they are reaching the same instruction
/// (hence all threads in the block are "aligned"). Also perform a fence before
/// and after the barrier according to \p Ordering. Note that the
/// fence might be part of the barrier if the target offers this.
[[gnu::noinline, omp::assume("ompx_aligned_barrier")]] void
threadsAligned(atomic::OrderingTy Ordering);
///}
} // namespace synchronize
// FIXME: NVPTX does not respect the memory scope argument.
namespace fence {
@ -261,6 +224,53 @@ static inline void system(atomic::OrderingTy Ordering) {
} // namespace fence
namespace synchronize {
/// Initialize the synchronization machinery. Must be called by all threads.
void init(bool IsSPMD);
/// Synchronize all threads in a warp identified by \p Mask.
static inline void warp(LaneMaskTy Mask) { __gpu_sync_lane(Mask); }
/// Synchronize all threads in a block and perform a fence before and after the
/// barrier according to \p Ordering. Note that the fence might be part of the
/// barrier.
static inline void threads(atomic::OrderingTy Ordering) {
#if defined(__NVPTX__)
__nvvm_barrier_sync(8);
#elif defined(__AMDGPU__)
if (Ordering != atomic::relaxed)
fence::team(Ordering == atomic::acq_rel ? atomic::release
: atomic::seq_cst);
__builtin_amdgcn_s_barrier();
if (Ordering != atomic::relaxed)
fence::team(Ordering == atomic::acq_rel ? atomic::acquire
: atomic::seq_cst);
#else
__gpu_sync_threads();
#endif
}
/// Synchronizing threads is allowed even if they all hit different instances of
/// `synchronize::threads()`. However, `synchronize::threadsAligned()` is more
/// restrictive in that it requires all threads to hit the same instance. The
/// noinline is removed by the openmp-opt pass and helps to preserve the
/// information till then.
///{
/// Synchronize all threads in a block, they are reaching the same instruction
/// (hence all threads in the block are "aligned"). Also perform a fence before
/// and after the barrier according to \p Ordering. Note that the
/// fence might be part of the barrier if the target offers this.
[[gnu::noinline, omp::assume("ompx_aligned_barrier")]] void
threadsAligned(atomic::OrderingTy Ordering);
///}
} // namespace synchronize
} // namespace ompx
#endif

8
openmp/device/src/Synchronization.cpp
View file

@ -33,7 +33,7 @@ namespace impl {
///{
#ifdef __AMDGPU__
[[clang::loader_uninitialized]] static Local<uint32_t> namedBarrierTracker;
[[clang::loader_uninitialized]] Local<uint32_t> namedBarrierTracker;
void namedBarrierInit() {
// Don't have global ctors, and shared memory is not zero init
@ -87,12 +87,6 @@ void namedBarrier() {
fence::team(atomic::release);
}
void syncWarp(__kmpc_impl_lanemask_t) {
// This is a no-op on current AMDGPU hardware but it is used by the optimizer
// to enforce convergent behaviour between control flow graphs.
__builtin_amdgcn_wave_barrier();
}
void syncThreadsAligned(atomic::OrderingTy Ordering) {
synchronize::threads(Ordering);
}