Skip to content

GPU: make the Metal atomics operate on plain counters #15944

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
ktf wants to merge 1 commit into
base: dev
Choose a base branch
from
+27 −10
Open

GPU: make the Metal atomics operate on plain counters #15944

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion GPU/Common/GPUCommonDefAPI.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -173,7 +173,7 @@
#define GPUdouble() float
#define GPUbarrier() threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)
#define GPUbarrierWarp() simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)
#define GPUAtomic(type) atomic<type> // atomic variable type
#define GPUAtomic(type) type // atomic variable type
#elif defined(__HIPCC__) //Defines for HIP
#define GPUd() __device__
#define GPUdDefault() __device__
Expand Down
27 changes: 22 additions & 5 deletions GPU/Common/GPUCommonMath.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -489,6 +489,23 @@ GPUhdi() constexpr int32_t GPUCommonMath::Abs<int32_t>(int32_t x)
return GPUCA_CHOICE(abs(x), abs(x), abs(x));
}

#ifdef __METAL__
// The counters these operate on are plain integers in the transferred structures,
// as they are for CUDA and HIP. MSL's atomic operations want metal::atomic, which
// has the same size and alignment; the overloads keep the address space, which a
// generic pointer would not carry into atomic_*_explicit.
template <class T>
GPUdi() threadgroup metal::atomic<T>* GPUCommonMathMetalAtomic(threadgroup T* p)
{
return reinterpret_cast<threadgroup metal::atomic<T>*>(p);
}
template <class T>
GPUdi() device metal::atomic<T>* GPUCommonMathMetalAtomic(T* p)
{
return (device metal::atomic<T>*)p;
}
#endif

template <class S, class T>
GPUdi() uint32_t GPUCommonMath::AtomicExchInternal(S* addr, T val)
{
Expand All @@ -499,7 +516,7 @@ GPUdi() uint32_t GPUCommonMath::AtomicExchInternal(S* addr, T val)
#elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__))
return ::atomicExch(addr, val);
#elif defined(GPUCA_GPUCODE) && defined(__METAL__)
return atomic_exchange_explicit(addr, val, memory_order_relaxed);
return atomic_exchange_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed);
#elif defined(WITH_OPENMP)
uint32_t old;
__atomic_exchange(addr, &val, &old, __ATOMIC_SEQ_CST);
Expand All @@ -519,7 +536,7 @@ GPUdi() bool GPUCommonMath::AtomicCASInternal(S* addr, T cmp, T val)
#elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__))
return ::atomicCAS(addr, cmp, val) == cmp;
#elif defined(GPUCA_GPUCODE) && defined(__METAL__)
return atomic_compare_exchange_weak_explicit(addr, &cmp, val, memory_order_relaxed, memory_order_relaxed);
return atomic_compare_exchange_weak_explicit(GPUCommonMathMetalAtomic(addr), &cmp, val, memory_order_relaxed, memory_order_relaxed);
#elif defined(WITH_OPENMP)
return __atomic_compare_exchange(addr, &cmp, &val, true, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST);
#else
Expand All @@ -537,7 +554,7 @@ GPUdi() uint32_t GPUCommonMath::AtomicAddInternal(S* addr, T val)
#elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__))
return ::atomicAdd(addr, val);
#elif defined(GPUCA_GPUCODE) && defined(__METAL__)
return atomic_fetch_add_explicit(addr, val, memory_order_relaxed);
return atomic_fetch_add_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed);
#elif defined(WITH_OPENMP)
return __atomic_add_fetch(addr, val, __ATOMIC_SEQ_CST) - val;
#else
Expand All @@ -555,7 +572,7 @@ GPUdi() void GPUCommonMath::AtomicMaxInternal(S* addr, T val)
#elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__))
::atomicMax(addr, val);
#elif defined(GPUCA_GPUCODE) && defined(__METAL__)
atomic_fetch_max_explicit(addr, val, memory_order_relaxed);
atomic_fetch_max_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed);
#else
S current;
while ((current = *(volatile S*)addr) < val && !AtomicCASInternal(addr, current, val)) {
Expand All @@ -573,7 +590,7 @@ GPUdi() void GPUCommonMath::AtomicMinInternal(S* addr, T val)
#elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__))
::atomicMin(addr, val);
#elif defined(GPUCA_GPUCODE) && defined(__METAL__)
atomic_fetch_min_explicit(addr, val, memory_order_relaxed);
atomic_fetch_min_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed);
#else
S current;
while ((current = *(volatile S*)addr) > val && !AtomicCASInternal(addr, current, val)) {
Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,12 +37,12 @@ using namespace o2::tpc::constants;
// ===========================================================================

template <>
GPUdii() void GPUTPCCFDecodeZS::Thread<GPUTPCCFDecodeZS::decodeZS>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, int32_t firstHBF, int32_t tpcTimeBinCut)
GPUdii() void GPUTPCCFDecodeZS::Thread<GPUTPCCFDecodeZS::decodeZS>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& clusterer, int32_t firstHBF, int32_t tpcTimeBinCut)
{
GPUTPCCFDecodeZS::decode(clusterer, smem, nBlocks, nThreads, iBlock, iThread, firstHBF, tpcTimeBinCut);
}

GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut)
GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUsharedref() GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut)
{
const uint32_t sector = clusterer.mISector;
#ifdef GPUCA_GPUCODE
Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ class GPUTPCCFDecodeZS : public GPUKernelTemplate
decodeZS,
};

static GPUd() void decode(GPUTPCClusterFinder& clusterer, GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut);
static GPUd() void decode(GPUTPCClusterFinder& clusterer, GPUsharedref() GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut);

typedef GPUTPCClusterFinder processorType;
GPUhdi() static processorType* Processor(GPUConstantMem& processors)
Expand All @@ -59,7 +59,7 @@ class GPUTPCCFDecodeZS : public GPUKernelTemplate
}

template <int32_t iKernel = defaultKernel, typename... Args>
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, Args... args);
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& clusterer, Args... args);
};

class GPUTPCCFDecodeZSLinkBase : public GPUKernelTemplate
Expand Down
Loading