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