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