CUDAUtilities/interface/AtomicPairCounter.h

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;  // in a "One to Many" association is the number of "One"
0025         uint32_t m;  // in a "One to Many" association is the total number of associations
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       // increment n by 1 and m by i.  return previous value
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   }  // namespace cuda
0056 }  // namespace cms
0057
0058 #endif  // HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h