xgboost
bitfield.h
Go to the documentation of this file.
1 
5 #ifndef XGBOOST_COMMON_BITFIELD_H_
6 #define XGBOOST_COMMON_BITFIELD_H_
7 
8 #include <algorithm>
9 #include <bitset>
10 #include <cinttypes>
11 #include <iostream>
12 #include <sstream>
13 #include <string>
14 #include <vector>
15 
16 #if defined(__CUDACC__)
17 #include <thrust/copy.h>
18 #include <thrust/device_ptr.h>
19 #include "device_helpers.cuh"
20 #endif // defined(__CUDACC__)
21 
22 #include "xgboost/span.h"
23 #include "common.h"
24 
25 namespace xgboost {
26 
27 #if defined(__CUDACC__)
28 using BitFieldAtomicType = unsigned long long; // NOLINT
29 
30 __forceinline__ __device__ BitFieldAtomicType AtomicOr(BitFieldAtomicType* address,
31  BitFieldAtomicType val) {
32  BitFieldAtomicType old = *address, assumed; // NOLINT
33  do {
34  assumed = old;
35  old = atomicCAS(address, assumed, val | assumed);
36  } while (assumed != old);
37 
38  return old;
39 }
40 
41 __forceinline__ __device__ BitFieldAtomicType AtomicAnd(BitFieldAtomicType* address,
42  BitFieldAtomicType val) {
43  BitFieldAtomicType old = *address, assumed; // NOLINT
44  do {
45  assumed = old;
46  old = atomicCAS(address, assumed, val & assumed);
47  } while (assumed != old);
48 
49  return old;
50 }
51 #endif // defined(__CUDACC__)
52 
58 template <typename VT, typename Direction, bool IsConst = false>
60  using value_type = std::conditional_t<IsConst, VT const, VT>; // NOLINT
61  using pointer = value_type*; // NOLINT
62 
63  static value_type constexpr kValueSize = sizeof(value_type) * 8;
64  static value_type constexpr kOne = 1; // force correct type.
65 
66  struct Pos {
67  std::remove_const_t<value_type> int_pos {0};
68  std::remove_const_t<value_type> bit_pos {0};
69  };
70 
71  private:
73  static_assert(!std::is_signed<VT>::value, "Must use unsiged type as underlying storage.");
74 
75  public:
77  Pos pos_v;
78  if (pos == 0) {
79  return pos_v;
80  }
81  pos_v.int_pos = pos / kValueSize;
82  pos_v.bit_pos = pos % kValueSize;
83  return pos_v;
84  }
85 
86  public:
87  BitFieldContainer() = default;
89  XGBOOST_DEVICE BitFieldContainer(BitFieldContainer const& other) : bits_{other.bits_} {}
90  BitFieldContainer &operator=(BitFieldContainer const &that) = default;
92 
95 
96  /*\brief Compute the size of needed memory allocation. The returned value is in terms
97  * of number of elements with `BitFieldContainer::value_type'.
98  */
99  XGBOOST_DEVICE static size_t ComputeStorageSize(size_t size) {
100  return common::DivRoundUp(size, kValueSize);
101  }
102 #if defined(__CUDA_ARCH__)
103  __device__ BitFieldContainer& operator|=(BitFieldContainer const& rhs) {
104  auto tid = blockIdx.x * blockDim.x + threadIdx.x;
105  size_t min_size = min(bits_.size(), rhs.bits_.size());
106  if (tid < min_size) {
107  bits_[tid] |= rhs.bits_[tid];
108  }
109  return *this;
110  }
111 #else
113  size_t min_size = std::min(bits_.size(), rhs.bits_.size());
114  for (size_t i = 0; i < min_size; ++i) {
115  bits_[i] |= rhs.bits_[i];
116  }
117  return *this;
118  }
119 #endif // #if defined(__CUDA_ARCH__)
120 
121 #if defined(__CUDA_ARCH__)
122  __device__ BitFieldContainer& operator&=(BitFieldContainer const& rhs) {
123  size_t min_size = min(bits_.size(), rhs.bits_.size());
124  auto tid = blockIdx.x * blockDim.x + threadIdx.x;
125  if (tid < min_size) {
126  bits_[tid] &= rhs.bits_[tid];
127  }
128  return *this;
129  }
130 #else
132  size_t min_size = std::min(bits_.size(), rhs.bits_.size());
133  for (size_t i = 0; i < min_size; ++i) {
134  bits_[i] &= rhs.bits_[i];
135  }
136  return *this;
137  }
138 #endif // defined(__CUDA_ARCH__)
139 
140 #if defined(__CUDA_ARCH__)
141  __device__ auto Set(value_type pos) {
142  Pos pos_v = Direction::Shift(ToBitPos(pos));
143  value_type& value = bits_[pos_v.int_pos];
144  value_type set_bit = kOne << pos_v.bit_pos;
145  using Type = typename dh::detail::AtomicDispatcher<sizeof(value_type)>::Type;
146  atomicOr(reinterpret_cast<Type *>(&value), set_bit);
147  }
148  __device__ void Clear(value_type pos) {
149  Pos pos_v = Direction::Shift(ToBitPos(pos));
150  value_type& value = bits_[pos_v.int_pos];
151  value_type clear_bit = ~(kOne << pos_v.bit_pos);
152  using Type = typename dh::detail::AtomicDispatcher<sizeof(value_type)>::Type;
153  atomicAnd(reinterpret_cast<Type *>(&value), clear_bit);
154  }
155 #else
156  void Set(value_type pos) {
157  Pos pos_v = Direction::Shift(ToBitPos(pos));
158  value_type& value = bits_[pos_v.int_pos];
159  value_type set_bit = kOne << pos_v.bit_pos;
160  value |= set_bit;
161  }
162  void Clear(value_type pos) {
163  Pos pos_v = Direction::Shift(ToBitPos(pos));
164  value_type& value = bits_[pos_v.int_pos];
165  value_type clear_bit = ~(kOne << pos_v.bit_pos);
166  value &= clear_bit;
167  }
168 #endif // defined(__CUDA_ARCH__)
169 
170  XGBOOST_DEVICE bool Check(Pos pos_v) const {
171  pos_v = Direction::Shift(pos_v);
172  SPAN_LT(pos_v.int_pos, bits_.size());
173  value_type const value = bits_[pos_v.int_pos];
174  value_type const test_bit = kOne << pos_v.bit_pos;
175  value_type result = test_bit & value;
176  return static_cast<bool>(result);
177  }
178  XGBOOST_DEVICE bool Check(value_type pos) const {
179  Pos pos_v = ToBitPos(pos);
180  return Check(pos_v);
181  }
182 
183  XGBOOST_DEVICE size_t Size() const { return kValueSize * bits_.size(); }
184 
185  XGBOOST_DEVICE pointer Data() const { return bits_.data(); }
186 
187  inline friend std::ostream &
189  os << "Bits " << "storage size: " << field.bits_.size() << "\n";
190  for (typename common::Span<value_type>::index_type i = 0; i < field.bits_.size(); ++i) {
191  std::bitset<BitFieldContainer<VT, Direction, IsConst>::kValueSize> bset(field.bits_[i]);
192  os << bset << "\n";
193  }
194  return os;
195  }
196 };
197 
198 // Bits start from left most bits (most significant bit).
199 template <typename VT, bool IsConst = false>
200 struct LBitsPolicy : public BitFieldContainer<VT, LBitsPolicy<VT, IsConst>, IsConst> {
202  using Pos = typename Container::Pos;
203  using value_type = typename Container::value_type; // NOLINT
204 
205  XGBOOST_DEVICE static Pos Shift(Pos pos) {
206  pos.bit_pos = Container::kValueSize - pos.bit_pos - Container::kOne;
207  return pos;
208  }
209 };
210 
211 // Bits start from right most bit (least significant bit) of each entry, but integer index
212 // is from left to right.
213 template <typename VT>
214 struct RBitsPolicy : public BitFieldContainer<VT, RBitsPolicy<VT>> {
216  using Pos = typename Container::Pos;
217  using value_type = typename Container::value_type; // NOLINT
218 
219  XGBOOST_DEVICE static Pos Shift(Pos pos) {
220  return pos;
221  }
222 };
223 
224 // Format: <Const><Direction>BitField<size of underlying type in bits>, underlying type
225 // must be unsigned.
228 
231 } // namespace xgboost
232 
233 #endif // XGBOOST_COMMON_BITFIELD_H_
xgboost::BitFieldContainer::kOne
static constexpr value_type kOne
Definition: bitfield.h:64
xgboost::BitFieldContainer::Data
XGBOOST_DEVICE pointer Data() const
Definition: bitfield.h:185
xgboost::BitFieldContainer::BitFieldContainer
BitFieldContainer()=default
xgboost::LBitsPolicy::value_type
typename Container::value_type value_type
Definition: bitfield.h:203
xgboost::BitFieldContainer< VT, LBitsPolicy< VT, false >, false >::pointer
value_type * pointer
Definition: bitfield.h:61
xgboost::BitFieldContainer::Clear
void Clear(value_type pos)
Definition: bitfield.h:162
xgboost::BitFieldContainer::ComputeStorageSize
static XGBOOST_DEVICE size_t ComputeStorageSize(size_t size)
Definition: bitfield.h:99
xgboost::BitFieldContainer
A non-owning type with auxiliary methods defined for manipulating bits.
Definition: bitfield.h:59
xgboost::common::Span::index_type
std::size_t index_type
Definition: span.h:421
xgboost::BitFieldContainer::Bits
XGBOOST_DEVICE common::Span< value_type > Bits()
Definition: bitfield.h:93
xgboost::BitFieldContainer::operator=
BitFieldContainer & operator=(BitFieldContainer const &that)=default
xgboost::BitFieldContainer::value_type
std::conditional_t< IsConst, VT const, VT > value_type
Definition: bitfield.h:60
xgboost::BitFieldContainer::Pos::int_pos
std::remove_const_t< value_type > int_pos
Definition: bitfield.h:67
xgboost::RBitsPolicy
Definition: bitfield.h:214
xgboost::BitFieldContainer::operator&=
BitFieldContainer & operator&=(BitFieldContainer const &rhs)
Definition: bitfield.h:131
span.h
xgboost::RBitsPolicy::Pos
typename Container::Pos Pos
Definition: bitfield.h:216
xgboost::BitFieldContainer::Check
XGBOOST_DEVICE bool Check(Pos pos_v) const
Definition: bitfield.h:170
xgboost::RBitsPolicy::Shift
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:219
xgboost::LBitsPolicy::Pos
typename Container::Pos Pos
Definition: bitfield.h:202
xgboost::LBitsPolicy
Definition: bitfield.h:200
xgboost::common::DivRoundUp
XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b)
Definition: common.h:85
xgboost::LBitsPolicy::Shift
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:205
xgboost::BitFieldContainer::Pos
Definition: bitfield.h:66
xgboost::BitFieldContainer::kValueSize
static constexpr value_type kValueSize
Definition: bitfield.h:63
SPAN_LT
#define SPAN_LT(lhs, rhs)
Definition: span.h:117
xgboost::BitFieldContainer::ToBitPos
static XGBOOST_DEVICE Pos ToBitPos(value_type pos)
Definition: bitfield.h:76
xgboost::BitFieldContainer::Set
void Set(value_type pos)
Definition: bitfield.h:156
common.h
Common utilities.
xgboost::RBitsPolicy::value_type
typename Container::value_type value_type
Definition: bitfield.h:217
xgboost::common::Span< value_type >
xgboost::BitFieldContainer::Pos::bit_pos
std::remove_const_t< value_type > bit_pos
Definition: bitfield.h:68
xgboost::BitFieldContainer::BitFieldContainer
XGBOOST_DEVICE BitFieldContainer(BitFieldContainer const &other)
Definition: bitfield.h:89
xgboost::common::Span::size
constexpr XGBOOST_DEVICE index_type size() const __span_noexcept
Definition: span.h:547
xgboost::common::Span::data
constexpr XGBOOST_DEVICE pointer data() const __span_noexcept
Definition: span.h:542
xgboost::BitFieldContainer::operator|=
BitFieldContainer & operator|=(BitFieldContainer const &rhs)
Definition: bitfield.h:112
xgboost::BitFieldContainer::Size
XGBOOST_DEVICE size_t Size() const
Definition: bitfield.h:183
xgboost::BitFieldContainer::BitFieldContainer
XGBOOST_DEVICE BitFieldContainer(common::Span< value_type > bits)
Definition: bitfield.h:88
xgboost::BitFieldContainer::Bits
XGBOOST_DEVICE common::Span< value_type const > Bits() const
Definition: bitfield.h:94
xgboost::BitFieldContainer::operator<<
friend std::ostream & operator<<(std::ostream &os, BitFieldContainer< VT, Direction, IsConst > field)
Definition: bitfield.h:188
XGBOOST_DEVICE
#define XGBOOST_DEVICE
Tag function as usable by device.
Definition: base.h:84
xgboost::BitFieldContainer::Check
XGBOOST_DEVICE bool Check(value_type pos) const
Definition: bitfield.h:178
xgboost
namespace of xgboost
Definition: base.h:110