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