File indexing completed on 2024-04-06 12:15:44
0001 #ifndef HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h
0002 #define HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h
0003
0004 #include <cstdint>
0005
0006 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
0007
0008 namespace cms {
0009 namespace cuda {
0010
0011 class AtomicPairCounter {
0012 public:
0013 using c_type = unsigned long long int;
0014
0015 AtomicPairCounter() {}
0016 AtomicPairCounter(c_type i) { counter.ac = i; }
0017
0018 __device__ __host__ AtomicPairCounter& operator=(c_type i) {
0019 counter.ac = i;
0020 return *this;
0021 }
0022
0023 struct Counters {
0024 uint32_t n;
0025 uint32_t m;
0026 };
0027
0028 union Atomic2 {
0029 Counters counters;
0030 c_type ac;
0031 };
0032
0033 static constexpr c_type incr = 1UL << 32;
0034
0035 __device__ __host__ Counters get() const { return counter.counters; }
0036
0037
0038 __host__ __device__ __forceinline__ Counters add(uint32_t i) {
0039 c_type c = i;
0040 c += incr;
0041 Atomic2 ret;
0042 #ifdef __CUDA_ARCH__
0043 ret.ac = atomicAdd(&counter.ac, c);
0044 #else
0045 ret.ac = counter.ac;
0046 counter.ac += c;
0047 #endif
0048 return ret.counters;
0049 }
0050
0051 private:
0052 Atomic2 counter;
0053 };
0054
0055 }
0056 }
0057
0058 #endif