[Openmp-commits] [openmp] r370210 - [libomptarget] Refactor syncthreads macro to inline function
Jon Chesterfield via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 28 07:22:35 PDT 2019
Author: jonchesterfield
Date: Wed Aug 28 07:22:35 2019
New Revision: 370210
URL: http://llvm.org/viewvc/llvm-project?rev=370210&view=rev
Log:
[libomptarget] Refactor syncthreads macro to inline function
Summary:
[libomptarget] Refactor syncthreads macro to inline function
See also abandoned D66846, split into this diff and others.
Rev 2 of D66855
Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers
Subscribers: openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D66861
Modified:
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu?rev=370210&r1=370209&r2=370210&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu Wed Aug 28 07:22:35 2019
@@ -553,8 +553,7 @@ EXTERN void __kmpc_get_team_static_memor
if (GetThreadIdInBlock() == 0) {
*frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
}
- // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
return;
}
ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
@@ -568,8 +567,7 @@ EXTERN void __kmpc_restore_team_static_m
if (is_shared)
return;
if (isSPMDExecutionMode) {
- // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
if (GetThreadIdInBlock() == 0) {
omptarget_nvptx_simpleMemoryManager.Release();
}
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu?rev=370210&r1=370209&r2=370210&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu Wed Aug 28 07:22:35 2019
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "omptarget-nvptx.h"
+#include "target_impl.h"
////////////////////////////////////////////////////////////////////////////////
// global data tables
@@ -106,7 +107,7 @@ EXTERN void __kmpc_spmd_kernel_init(int
}
if (!RequiresOMPRuntime) {
// Runtime is not required - exit.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
return;
}
@@ -125,8 +126,7 @@ EXTERN void __kmpc_spmd_kernel_init(int
// init team context
currTeamDescr.InitTeamDescr();
}
- // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
@@ -168,8 +168,7 @@ EXTERN void __kmpc_spmd_kernel_deinit_v2
if (!RequiresOMPRuntime)
return;
- // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
int threadId = GetThreadIdInBlock();
if (threadId == 0) {
// Enqueue omp state object for use by another team.
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h?rev=370210&r1=370209&r2=370210&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h Wed Aug 28 07:22:35 2019
@@ -56,14 +56,6 @@
#define __ACTIVEMASK() __ballot(1)
#endif // CUDA_VERSION
-#define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
-// Use original __syncthreads if compiled by nvcc or clang >= 9.0.
-#if !defined(__clang__) || __clang_major__ >= 9
-#define __SYNCTHREADS() __syncthreads()
-#else
-#define __SYNCTHREADS() __SYNCTHREADS_N(0)
-#endif
-
// arguments needed for L0 parallelism only.
class omptarget_nvptx_SharedArgs {
public:
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu?rev=370210&r1=370209&r2=370210&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu Wed Aug 28 07:22:35 2019
@@ -75,8 +75,7 @@ EXTERN void __kmpc_barrier(kmp_Ident *lo
// parallel region and that all worker threads participate.
EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {
PRINT0(LD_SYNC, "call kmpc_barrier_simple_spmd\n");
- // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
- __SYNCTHREADS();
+ __kmpc_impl_syncthreads();
PRINT0(LD_SYNC, "completed kmpc_barrier_simple_spmd\n");
}
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h?rev=370210&r1=370209&r2=370210&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h Wed Aug 28 07:22:35 2019
@@ -63,6 +63,15 @@ INLINE int32_t __kmpc_impl_shfl_down_syn
#endif // CUDA_VERSION
}
+INLINE void __kmpc_impl_syncthreads() {
+ // Use original __syncthreads if compiled by nvcc or clang >= 9.0.
+#if !defined(__clang__) || __clang_major__ >= 9
+ __syncthreads();
+#else
+ asm volatile("bar.sync %0;" : : "r"(0) : "memory");
+#endif // __clang__
+}
+
INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
#if CUDA_VERSION >= 9000
__syncwarp(Mask);
More information about the Openmp-commits
mailing list