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>
72 static_assert(!std::is_signed<VT>::value,
"Must use unsiged type as underlying storage.");
98 #if defined(__CUDA_ARCH__) 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];
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];
115 #endif // #if defined(__CUDA_ARCH__) 117 #if defined(__CUDA_ARCH__) 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];
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];
134 #endif // defined(__CUDA_ARCH__) 136 #if defined(__CUDA_ARCH__) 138 Pos pos_v = Direction::Shift(ToBitPos(pos));
141 static_assert(
sizeof(BitFieldAtomicType) ==
sizeof(
value_type),
"");
142 AtomicOr(reinterpret_cast<BitFieldAtomicType*>(&value), set_bit);
145 Pos pos_v = Direction::Shift(ToBitPos(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);
153 Pos pos_v = Direction::Shift(ToBitPos(pos));
159 Pos pos_v = Direction::Shift(ToBitPos(pos));
161 value_type clear_bit = ~(kOne << pos_v.bit_pos);
164 #endif // defined(__CUDA_ARCH__) 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;
171 return static_cast<bool>(result);
174 Pos pos_v = ToBitPos(pos);
182 friend std::ostream& operator<<(std::ostream& os, BitFieldContainer<VT, Direction> field) {
183 os <<
"Bits " <<
"storage size: " << field.bits_.size() <<
"\n";
185 std::bitset<BitFieldContainer<VT, Direction>::kValueSize> bset(field.bits_[i]);
193 template <
typename VT>
196 using Pos =
typename Container::Pos;
200 pos.bit_pos = Container::kValueSize - pos.bit_pos - Container::kOne;
207 template <
typename VT>
210 using Pos =
typename Container::Pos;
222 #if defined(__CUDACC__) 224 template <
typename V,
typename D>
226 std::cout <<
"Bits: " << name << std::endl;
227 std::vector<typename BitFieldContainer<V, D>::value_type> h_field_bits(field.bits_.size());
230 field.bits_.data() + field.bits_.size()),
231 h_field_bits.data());
233 h_field.bits_ = {h_field_bits.
data(), h_field_bits.data() + h_field_bits.size()};
234 std::cout << h_field;
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()),
243 for (
auto v : h_list) {
244 std::cout << v <<
", ";
246 std::cout << std::endl;
249 #endif // defined(__CUDACC__) 252 #endif // XGBOOST_COMMON_BITFIELD_H_ XGBOOST_DEVICE constexpr index_type size() const __span_noexcept
Definition: span.h:531
Definition: bitfield.h:208
std::size_t index_type
Definition: span.h:405
XGBOOST_DEVICE size_t Size() const
Definition: bitfield.h:178
VT value_type
Definition: bitfield.h:59
common::Span< value_type const > Bits() const
Definition: bitfield.h:90
XGBOOST_DEVICE bool Check(Pos pos_v) const
Definition: bitfield.h:166
BitFieldContainer & operator &=(BitFieldContainer const &rhs)
Definition: bitfield.h:127
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:213
static value_type constexpr kOne
Definition: bitfield.h:63
static size_t ComputeStorageSize(size_t size)
Definition: bitfield.h:95
value_type bit_pos
Definition: bitfield.h:67
XGBOOST_DEVICE BitFieldContainer(BitFieldContainer const &other)
Definition: bitfield.h:87
XGBOOST_DEVICE constexpr pointer data() const __span_noexcept
Definition: span.h:526
XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b)
Definition: common.h:81
XGBOOST_DEVICE pointer Data() const
Definition: bitfield.h:180
static value_type constexpr kValueSize
Definition: bitfield.h:62
typename Container::value_type value_type
Definition: bitfield.h:211
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:199
typename Container::Pos Pos
Definition: bitfield.h:210
Definition: bitfield.h:65
#define XGBOOST_DEVICE
Tag function as usable by device.
Definition: base.h:84
typename Container::Pos Pos
Definition: bitfield.h:196
void Clear(value_type pos)
Definition: bitfield.h:158
namespace of xgboost
Definition: base.h:102
value_type int_pos
Definition: bitfield.h:66
typename Container::value_type value_type
Definition: bitfield.h:197
XGBOOST_DEVICE BitFieldContainer(common::Span< value_type > bits)
Definition: bitfield.h:86
void Set(value_type pos)
Definition: bitfield.h:152
common::Span< value_type > Bits()
Definition: bitfield.h:89
value_type * pointer
Definition: bitfield.h:60
Definition: bitfield.h:194
BitFieldContainer & operator|=(BitFieldContainer const &rhs)
Definition: bitfield.h:108
XGBOOST_DEVICE bool Check(value_type pos) const
Definition: bitfield.h:173
BitFieldContainer()=default
A non-owning type with auxiliary methods defined for manipulating bits.
Definition: bitfield.h:58