Line data    Source code

       1             : #pragma once
       2             : 
       3             : #include <cinttypes>
       4             : #include <type_traits>
       5             : 
       6             : /* define guards for __host__ and __device__ tags used by cuda */
       7             : #ifndef __host__
       8             : #define __host__
       9             : #endif
      10             : #ifndef __device__
      11             : #define __device__
      12             : #endif
      13             : 
      14             : namespace elsa::mr::detail
      15             : {
      16             :     /*
      17             :      *  Wraps a pointer and offers 'get' function to retrieve it.
      18             :      *  Pointer is expected to be accessible on hosts and devices.
      19             :      */
      20             :     template <class Type>
      21             :     class ContPointer
      22             :     {
      23             :     public:
      24             :         using self_type = ContPointer<Type>;
      25             :         using value_type = std::remove_cv_t<Type>;
      26             :         using size_type = size_t;
      27             :         using difference_type = std::ptrdiff_t;
      28             : 
      29             :         using raw_pointer = Type*;
      30             :         using pointer = Type*;
      31             :         using reference = Type&;
      32             : 
      33             :     private:
      34             :         pointer _where = 0;
      35             : 
      36             :     public:
      37             :         __host__ __device__ ContPointer() {}
      38         125 :         __host__ __device__ ContPointer(pointer w) : _where(w) {}
      39             :         __host__ __device__ ContPointer(self_type&& p) noexcept : _where(p._where) {}
      40             :         __host__ __device__ ContPointer(const ContPointer<std::remove_const_t<Type>>& p)
      41             :             : _where(p.get())
      42          30 :         {
      43          30 :         }
      44             :         __host__ __device__ ContPointer(const ContPointer<std::add_const_t<Type>>& p)
      45             :             : _where(p.get())
      46          15 :         {
      47          15 :             static_assert(std::is_const<Type>::value,
      48          15 :                           "Const pointer cannot be converted to a normal pointer");
      49          15 :         }
      50             : 
      51             :     public:
      52          60 :         __host__ __device__ raw_pointer get() const { return _where; }
      53             : 
      54             :     public:
      55             :         __host__ __device__ self_type& operator=(const self_type& p)
      56             :         {
      57             :             _where = p._where;
      58             :             return *this;
      59             :         }
      60             :         __host__ __device__ self_type& operator=(self_type&& p) noexcept
      61             :         {
      62             :             _where = p._where;
      63             :             return *this;
      64             :         }
      65          55 :         __host__ __device__ bool operator==(const self_type& p) const { return _where == p._where; }
      66           5 :         __host__ __device__ bool operator!=(const self_type& p) const { return !(*this == p); }
      67             :         __host__ __device__ reference operator*() const { return *_where; }
      68             :         __host__ __device__ pointer operator->() const { return _where; }
      69             :         __host__ __device__ self_type& operator++()
      70             :         {
      71             :             ++_where;
      72             :             return *this;
      73             :         };
      74             :         __host__ __device__ self_type operator++(int)
      75             :         {
      76             :             self_type out(_where);
      77             :             ++_where;
      78             :             return out;
      79             :         }
      80             :         __host__ __device__ self_type& operator--()
      81             :         {
      82             :             --_where;
      83             :             return *this;
      84             :         };
      85             :         __host__ __device__ self_type operator--(int)
      86             :         {
      87             :             self_type out(_where);
      88             :             --_where;
      89             :             return out;
      90             :         }
      91             : 
      92             :         __host__ __device__ self_type& operator+=(difference_type d)
      93             :         {
      94             :             _where += d;
      95             :             return *this;
      96             :         }
      97             :         __host__ __device__ self_type& operator-=(difference_type d)
      98             :         {
      99             :             _where -= d;
     100             :             return *this;
     101             :         }
     102             :         __host__ __device__ self_type operator+(difference_type d) const
     103             :         {
     104             :             return self_type(_where + d);
     105             :         }
     106             :         __host__ __device__ self_type operator-(difference_type d) const
     107             :         {
     108             :             return self_type(_where - d);
     109             :         }
     110             :         __host__ __device__ difference_type operator-(const self_type& p) const
     111             :         {
     112             :             return _where - p._where;
     113             :         }
     114             :         __host__ __device__ reference operator[](size_type i) const { return _where[i]; }
     115             :         __host__ __device__ bool operator<(const self_type& p) const { return _where < p._where; }
     116             :         __host__ __device__ bool operator<=(const self_type& p) const { return _where <= p._where; }
     117             :         __host__ __device__ bool operator>=(const self_type& p) const { return !(*this < p); }
     118             :         __host__ __device__ bool operator>(const self_type& p) const { return !(*this <= p); }
     119             :     };
     120             : 
     121             :     /*
     122             :      *  Wraps a pointer to a contiguous (random-access) range of values.
     123             :      *  Pointer is expected to be accessible on hosts and devices.
     124             :      */
     125             :     template <class Type>
     126             :     class ContIterator
     127             :     {
     128             :     public:
     129             :         using self_type = ContIterator<Type>;
     130             :         using value_type = std::remove_cv_t<Type>;
     131             :         using iterator_category = std::random_access_iterator_tag;
     132             :         using difference_type = std::ptrdiff_t;
     133             :         using size_type = size_t;
     134             :         using pointer = Type*;
     135             :         using reference = Type&;
     136             : 
     137             :     private:
     138             :         pointer _where = 0;
     139             : 
     140             :     public:
     141             :         __host__ __device__ ContIterator() {}
     142        1291 :         __host__ __device__ ContIterator(pointer w) : _where(w) {}
     143          10 :         __host__ __device__ ContIterator(self_type&& p) noexcept : _where(p._where) {}
     144             :         __host__ __device__ ContIterator(const ContIterator<std::remove_const_t<Type>>& p)
     145             :             : _where(p.base())
     146         649 :         {
     147         649 :         }
     148             :         __host__ __device__ ContIterator(const ContIterator<std::add_const_t<Type>>& p)
     149             :             : _where(p.base())
     150        1043 :         {
     151        1043 :             static_assert(std::is_const<Type>::value,
     152        1043 :                           "Const iterator cannot be converted to a normal iterator");
     153        1043 :         }
     154             : 
     155             :     public:
     156             :         __host__ __device__ self_type& operator=(const self_type& p)
     157             :         {
     158             :             _where = p._where;
     159             :             return *this;
     160             :         }
     161             :         __host__ __device__ self_type& operator=(self_type&& p) noexcept
     162             :         {
     163             :             _where = p._where;
     164             :             return *this;
     165             :         }
     166          90 :         __host__ __device__ bool operator==(const self_type& p) const { return _where == p._where; }
     167             :         __host__ __device__ bool operator!=(const self_type& p) const { return !(*this == p); }
     168       17797 :         __host__ __device__ reference operator*() const { return *_where; }
     169             :         __host__ __device__ pointer operator->() const { return _where; }
     170             :         __host__ __device__ self_type& operator++()
     171       16405 :         {
     172       16405 :             ++_where;
     173       16405 :             return *this;
     174       16405 :         };
     175             :         __host__ __device__ self_type operator++(int)
     176             :         {
     177             :             self_type out(_where);
     178             :             ++_where;
     179             :             return out;
     180             :         }
     181             :         __host__ __device__ self_type& operator--()
     182        2709 :         {
     183        2709 :             --_where;
     184        2709 :             return *this;
     185        2709 :         };
     186             :         __host__ __device__ self_type operator--(int)
     187             :         {
     188             :             self_type out(_where);
     189             :             --_where;
     190             :             return out;
     191             :         }
     192             : 
     193             :         __host__ __device__ self_type& operator+=(difference_type d)
     194          10 :         {
     195          10 :             _where += d;
     196          10 :             return *this;
     197          10 :         }
     198             :         __host__ __device__ self_type& operator-=(difference_type d)
     199             :         {
     200             :             _where -= d;
     201             :             return *this;
     202             :         }
     203             :         __host__ __device__ self_type operator+(difference_type d) const
     204         195 :         {
     205         195 :             return self_type(_where + d);
     206         195 :         }
     207             :         __host__ __device__ self_type operator-(difference_type d) const
     208             :         {
     209             :             return self_type(_where - d);
     210             :         }
     211             :         __host__ __device__ difference_type operator-(const self_type& p) const
     212         170 :         {
     213         170 :             return _where - p._where;
     214         170 :         }
     215             :         __host__ __device__ reference operator[](size_type i) const { return _where[i]; }
     216             :         __host__ __device__ bool operator<(const self_type& p) const { return _where < p._where; }
     217             :         __host__ __device__ bool operator<=(const self_type& p) const { return _where <= p._where; }
     218             :         __host__ __device__ bool operator>=(const self_type& p) const { return !(*this < p); }
     219             :         __host__ __device__ bool operator>(const self_type& p) const { return !(*this <= p); }
     220        1692 :         __host__ __device__ pointer base() const { return _where; }
     221             :     };
     222             : } // namespace elsa::mr::detail