gpuClustering::pixelStatus Namespace Reference

Enumerations
enum	Status : uint32_t { kEmpty = 0x00, kFound = 0x01, kDuplicate = 0x03 }

Functions
static __device__ constexpr uint32_t	getIndex (uint16_t x, uint16_t y)
__device__ constexpr uint32_t	getShift (uint16_t x, uint16_t y)
__device__ constexpr Status	getStatus (uint32_t const *__restrict__ status, uint16_t x, uint16_t y)
__device__ constexpr bool	isDuplicate (uint32_t const *__restrict__ status, uint16_t x, uint16_t y)
__device__ constexpr void	promote (uint32_t *__restrict__ status, const uint16_t x, const uint16_t y)

Variables
constexpr uint32_t	bits = 2
constexpr uint32_t	mask = (0x01 << bits) - 1
constexpr uint32_t	size = pixelSizeX * pixelSizeY / valuesPerWord
constexpr uint32_t	valuesPerWord = sizeof(uint32_t) * 8 / bits

Enumeration Type Documentation

Status

enum gpuClustering::pixelStatus::Status : uint32_t

Enumerator
kEmpty
kFound
kDuplicate

Definition at line 23 of file gpuClustering.h.

: uint32_t { kEmpty = 0x00, kFound = 0x01, kDuplicate = 0x03 };

Function Documentation

getIndex()

static __device__ constexpr uint32_t gpuClustering::pixelStatus::getIndex ( uint16_t  x, uint16_t  y  )

inlinestatic

Definition at line 30 of file gpuClustering.h.

References gpuClustering::pixelSizeX, valuesPerWord, and x.

Referenced by getStatus(), and promote().

                                                                                 {
      return (pixelSizeX * y + x) / valuesPerWord;
    }

getShift()

__device__ constexpr uint32_t gpuClustering::pixelStatus::getShift ( uint16_t  x, uint16_t  y  )

inline

Definition at line 34 of file gpuClustering.h.

References valuesPerWord, and x.

Referenced by getStatus(), and promote().

{ return (x % valuesPerWord) * 2; }

getStatus()

__device__ constexpr Status gpuClustering::pixelStatus::getStatus ( uint32_t const *__restrict__  status, uint16_t  x, uint16_t  y  )

inline

Definition at line 36 of file gpuClustering.h.

References getIndex(), getShift(), mask, edm::shift, gpuClustering::status, and x.

Referenced by isDuplicate().

                                                                                                              {
      uint32_t index = getIndex(x, y);
      uint32_t shift = getShift(x, y);
      return Status{(status[index] >> shift) & mask};
    }

isDuplicate()

__device__ constexpr bool gpuClustering::pixelStatus::isDuplicate ( uint32_t const *__restrict__  status, uint16_t  x, uint16_t  y  )

inline

Definition at line 42 of file gpuClustering.h.

References getStatus(), kDuplicate, gpuClustering::status, and x.

                                                                                                              {
      return getStatus(status, x, y) == kDuplicate;
    }

promote()

__device__ constexpr void gpuClustering::pixelStatus::promote ( uint32_t *__restrict__  status, const uint16_t  x, const uint16_t  y  )

inline

Definition at line 46 of file gpuClustering.h.

References cms::cudacompat::atomicCAS(), getIndex(), getShift(), kDuplicate, kEmpty, kFound, mask, edm::shift, gpuClustering::status, and x.

                                                                                                                {
      uint32_t index = getIndex(x, y);
      uint32_t shift = getShift(x, y);
      uint32_t old_word = status[index];
      uint32_t expected = old_word;
      do {
        expected = old_word;
        Status old_status{(old_word >> shift) & mask};
        if (kDuplicate == old_status) {
          // nothing to do
          return;
        }
        Status new_status = (kEmpty == old_status) ? kFound : kDuplicate;
        uint32_t new_word = old_word | (static_cast<uint32_t>(new_status) << shift);
        old_word = atomicCAS(&status[index], expected, new_word);
      } while (expected != old_word);
    }

Variable Documentation

bits

constexpr uint32_t gpuClustering::pixelStatus::bits = 2

Definition at line 25 of file gpuClustering.h.

mask

constexpr uint32_t gpuClustering::pixelStatus::mask = (0x01 << bits) - 1

Definition at line 26 of file gpuClustering.h.

Referenced by getStatus(), and promote().

size

constexpr uint32_t gpuClustering::pixelStatus::size = pixelSizeX * pixelSizeY / valuesPerWord

Definition at line 28 of file gpuClustering.h.

valuesPerWord

constexpr uint32_t gpuClustering::pixelStatus::valuesPerWord = sizeof(uint32_t) * 8 / bits

Definition at line 27 of file gpuClustering.h.

Referenced by getIndex(), and getShift().