Go to the documentation of this file.
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 #include "device_helpers.cuh"
20 #endif // defined(__CUDACC__)
27 #if defined(__CUDACC__)
28 using BitFieldAtomicType =
unsigned long long;
30 __forceinline__ __device__ BitFieldAtomicType AtomicOr(BitFieldAtomicType* address,
31 BitFieldAtomicType val) {
32 BitFieldAtomicType old = *address, assumed;
35 old = atomicCAS(address, assumed, val | assumed);
36 }
while (assumed != old);
41 __forceinline__ __device__ BitFieldAtomicType AtomicAnd(BitFieldAtomicType* address,
42 BitFieldAtomicType val) {
43 BitFieldAtomicType old = *address, assumed;
46 old = atomicCAS(address, assumed, val & assumed);
47 }
while (assumed != old);
51 #endif // defined(__CUDACC__)
58 template <
typename VT,
typename Direction,
bool IsConst = false>
60 using value_type = std::conditional_t<IsConst, VT const, VT>;
67 std::remove_const_t<value_type>
int_pos {0};
68 std::remove_const_t<value_type>
bit_pos {0};
73 static_assert(!std::is_signed<VT>::value,
"Must use unsiged type as underlying storage.");
102 #if defined(__CUDA_ARCH__)
104 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
105 size_t min_size = min(bits_.
size(), rhs.bits_.
size());
106 if (tid < min_size) {
107 bits_[tid] |= rhs.bits_[tid];
113 size_t min_size = std::min(bits_.
size(), rhs.bits_.
size());
114 for (
size_t i = 0; i < min_size; ++i) {
115 bits_[i] |= rhs.bits_[i];
119 #endif // #if defined(__CUDA_ARCH__)
121 #if defined(__CUDA_ARCH__)
123 size_t min_size = min(bits_.
size(), rhs.bits_.
size());
124 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
125 if (tid < min_size) {
126 bits_[tid] &= rhs.bits_[tid];
132 size_t min_size = std::min(bits_.
size(), rhs.bits_.
size());
133 for (
size_t i = 0; i < min_size; ++i) {
134 bits_[i] &= rhs.bits_[i];
138 #endif // defined(__CUDA_ARCH__)
140 #if defined(__CUDA_ARCH__)
142 Pos pos_v = Direction::Shift(
ToBitPos(pos));
145 using Type =
typename dh::detail::AtomicDispatcher<
sizeof(
value_type)>::Type;
146 atomicOr(
reinterpret_cast<Type *
>(&value), set_bit);
149 Pos pos_v = Direction::Shift(
ToBitPos(pos));
152 using Type =
typename dh::detail::AtomicDispatcher<
sizeof(
value_type)>::Type;
153 atomicAnd(
reinterpret_cast<Type *
>(&value), clear_bit);
157 Pos pos_v = Direction::Shift(
ToBitPos(pos));
163 Pos pos_v = Direction::Shift(
ToBitPos(pos));
168 #endif // defined(__CUDA_ARCH__)
171 pos_v = Direction::Shift(pos_v);
173 value_type const value = bits_[pos_v.int_pos];
176 return static_cast<bool>(result);
187 inline friend std::ostream &
189 os <<
"Bits " <<
"storage size: " << field.bits_.size() <<
"\n";
191 std::bitset<BitFieldContainer<VT, Direction, IsConst>::kValueSize> bset(field.bits_[i]);
199 template <
typename VT,
bool IsConst = false>
213 template <
typename VT>
216 using Pos =
typename Container::Pos;
233 #endif // XGBOOST_COMMON_BITFIELD_H_
static constexpr value_type kOne
Definition: bitfield.h:64
XGBOOST_DEVICE pointer Data() const
Definition: bitfield.h:185
BitFieldContainer()=default
typename Container::value_type value_type
Definition: bitfield.h:203
value_type * pointer
Definition: bitfield.h:61
void Clear(value_type pos)
Definition: bitfield.h:162
static XGBOOST_DEVICE size_t ComputeStorageSize(size_t size)
Definition: bitfield.h:99
A non-owning type with auxiliary methods defined for manipulating bits.
Definition: bitfield.h:59
std::size_t index_type
Definition: span.h:421
XGBOOST_DEVICE common::Span< value_type > Bits()
Definition: bitfield.h:93
BitFieldContainer & operator=(BitFieldContainer const &that)=default
std::conditional_t< IsConst, VT const, VT > value_type
Definition: bitfield.h:60
std::remove_const_t< value_type > int_pos
Definition: bitfield.h:67
Definition: bitfield.h:214
BitFieldContainer & operator&=(BitFieldContainer const &rhs)
Definition: bitfield.h:131
typename Container::Pos Pos
Definition: bitfield.h:216
XGBOOST_DEVICE bool Check(Pos pos_v) const
Definition: bitfield.h:170
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:219
typename Container::Pos Pos
Definition: bitfield.h:202
Definition: bitfield.h:200
XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b)
Definition: common.h:85
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:205
Definition: bitfield.h:66
static constexpr value_type kValueSize
Definition: bitfield.h:63
#define SPAN_LT(lhs, rhs)
Definition: span.h:117
static XGBOOST_DEVICE Pos ToBitPos(value_type pos)
Definition: bitfield.h:76
void Set(value_type pos)
Definition: bitfield.h:156
typename Container::value_type value_type
Definition: bitfield.h:217
std::remove_const_t< value_type > bit_pos
Definition: bitfield.h:68
XGBOOST_DEVICE BitFieldContainer(BitFieldContainer const &other)
Definition: bitfield.h:89
constexpr XGBOOST_DEVICE index_type size() const __span_noexcept
Definition: span.h:547
constexpr XGBOOST_DEVICE pointer data() const __span_noexcept
Definition: span.h:542
BitFieldContainer & operator|=(BitFieldContainer const &rhs)
Definition: bitfield.h:112
XGBOOST_DEVICE size_t Size() const
Definition: bitfield.h:183
XGBOOST_DEVICE BitFieldContainer(common::Span< value_type > bits)
Definition: bitfield.h:88
XGBOOST_DEVICE common::Span< value_type const > Bits() const
Definition: bitfield.h:94
friend std::ostream & operator<<(std::ostream &os, BitFieldContainer< VT, Direction, IsConst > field)
Definition: bitfield.h:188
#define XGBOOST_DEVICE
Tag function as usable by device.
Definition: base.h:84
XGBOOST_DEVICE bool Check(value_type pos) const
Definition: bitfield.h:178
namespace of xgboost
Definition: base.h:110