AtomicPairCounter.h

Go to the documentation of this file.

#ifndef HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h
#define HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h

#include <cstdint>

#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"

namespace cms {
  namespace cuda {

    class AtomicPairCounter {
    public:
      using c_type = unsigned long long int;

      AtomicPairCounter() {}
      AtomicPairCounter(c_type i) { counter.ac = i; }

      __device__ __host__ AtomicPairCounter& operator=(c_type i) {
        counter.ac = i;
        return *this;
      }

      struct Counters {
        uint32_t n;  // in a "One to Many" association is the number of "One"
        uint32_t m;  // in a "One to Many" association is the total number of associations
      };

      union Atomic2 {
        Counters counters;
        c_type ac;
      };

      static constexpr c_type incr = 1UL << 32;

      __device__ __host__ Counters get() const { return counter.counters; }

      // increment n by 1 and m by i.  return previous value
      __host__ __device__ __forceinline__ Counters add(uint32_t i) {
        c_type c = i;
        c += incr;
        Atomic2 ret;
#ifdef __CUDA_ARCH__
        ret.ac = atomicAdd(&counter.ac, c);
#else
        ret.ac = counter.ac;
        counter.ac += c;
#endif
        return ret.counters;
      }

    private:
      Atomic2 counter;
    };

  }  // namespace cuda
}  // namespace cms

#endif  // HeterogeneousCore_CUDAUtilities_interface_AtomicPairCounter_h