5 #ifndef XGBOOST_COMMON_BITFIELD_H_ 6 #define XGBOOST_COMMON_BITFIELD_H_ 16 #if defined(__CUDACC__) 17 #include <thrust/copy.h> 18 #include <thrust/device_ptr.h> 19 #endif // defined(__CUDACC__) 26 #if defined(__CUDACC__) 27 using BitFieldAtomicType =
unsigned long long;
29 __forceinline__ __device__ BitFieldAtomicType AtomicOr(BitFieldAtomicType* address,
30 BitFieldAtomicType val) {
31 BitFieldAtomicType old = *address, assumed;
34 old = atomicCAS(address, assumed, val | assumed);
35 }
while (assumed != old);
40 __forceinline__ __device__ BitFieldAtomicType AtomicAnd(BitFieldAtomicType* address,
41 BitFieldAtomicType val) {
42 BitFieldAtomicType old = *address, assumed;
45 old = atomicCAS(address, assumed, val & assumed);
46 }
while (assumed != old);
50 #endif // defined(__CUDACC__) 57 template <
typename VT,
typename Direction>
71 static_assert(!std::is_signed<VT>::value,
"Must use unsiged type as underlying storage.");
94 #if defined(__CUDA_ARCH__) 96 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
99 bits_[tid] |= rhs.
bits_[tid];
106 for (
size_t i = 0; i < min_size; ++i) {
107 bits_[i] |= rhs.
bits_[i];
111 #endif // #if defined(__CUDA_ARCH__) 113 #if defined(__CUDA_ARCH__) 116 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
117 if (tid < min_size) {
118 bits_[tid] &= rhs.
bits_[tid];
125 for (
size_t i = 0; i < min_size; ++i) {
126 bits_[i] &= rhs.
bits_[i];
130 #endif // defined(__CUDA_ARCH__) 132 #if defined(__CUDA_ARCH__) 134 Pos pos_v = Direction::Shift(
ToBitPos(pos));
137 static_assert(
sizeof(BitFieldAtomicType) ==
sizeof(
value_type),
"");
138 AtomicOr(reinterpret_cast<BitFieldAtomicType*>(&value), set_bit);
141 Pos pos_v = Direction::Shift(
ToBitPos(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);
149 Pos pos_v = Direction::Shift(
ToBitPos(pos));
155 Pos pos_v = Direction::Shift(
ToBitPos(pos));
157 value_type clear_bit = ~(kOne << pos_v.bit_pos);
160 #endif // defined(__CUDA_ARCH__) 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;
167 return static_cast<bool>(result);
178 friend std::ostream& operator<<(std::ostream& os, BitFieldContainer<VT, Direction> field) {
179 os <<
"Bits " <<
"storage size: " << field.bits_.size() <<
"\n";
181 std::bitset<BitFieldContainer<VT, Direction>::kValueSize> bset(field.bits_[i]);
189 template <
typename VT>
192 using Pos =
typename Container::Pos;
196 pos.bit_pos = Container::kValueSize - pos.bit_pos - Container::kOne;
203 template <
typename VT>
206 using Pos =
typename Container::Pos;
218 #if defined(__CUDACC__) 220 template <
typename V,
typename D>
222 std::cout <<
"Bits: " << name << std::endl;
223 std::vector<typename BitFieldContainer<V, D>::value_type> h_field_bits(field.
bits_.size());
227 h_field_bits.data());
229 h_field.
bits_ = {h_field_bits.
data(), h_field_bits.data() + h_field_bits.size()};
230 std::cout << h_field;
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()),
239 for (
auto v : h_list) {
240 std::cout << v <<
", ";
242 std::cout << std::endl;
245 #endif // defined(__CUDACC__) 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
VT value_type
Definition: bitfield.h:59
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
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
BitFieldContainer()=default
A non-owning type with auxiliary methods defined for manipulating bits.
Definition: bitfield.h:58