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.");
100 #if defined(__CUDA_ARCH__)
102 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
103 size_t min_size = min(bits_.
size(), rhs.bits_.
size());
104 if (tid < min_size) {
105 bits_[tid] |= rhs.bits_[tid];
111 size_t min_size = std::min(bits_.
size(), rhs.bits_.
size());
112 for (
size_t i = 0; i < min_size; ++i) {
113 bits_[i] |= rhs.bits_[i];
117 #endif // #if defined(__CUDA_ARCH__)
119 #if defined(__CUDA_ARCH__)
121 size_t min_size = min(bits_.
size(), rhs.bits_.
size());
122 auto tid = blockIdx.x * blockDim.x + threadIdx.x;
123 if (tid < min_size) {
124 bits_[tid] &= rhs.bits_[tid];
130 size_t min_size = std::min(bits_.
size(), rhs.bits_.
size());
131 for (
size_t i = 0; i < min_size; ++i) {
132 bits_[i] &= rhs.bits_[i];
136 #endif // defined(__CUDA_ARCH__)
138 #if defined(__CUDA_ARCH__)
140 Pos pos_v = Direction::Shift(
ToBitPos(pos));
143 using Type =
typename dh::detail::AtomicDispatcher<
sizeof(
value_type)>::Type;
144 atomicOr(
reinterpret_cast<Type *
>(&value), set_bit);
147 Pos pos_v = Direction::Shift(
ToBitPos(pos));
150 using Type =
typename dh::detail::AtomicDispatcher<
sizeof(
value_type)>::Type;
151 atomicAnd(
reinterpret_cast<Type *
>(&value), clear_bit);
155 Pos pos_v = Direction::Shift(
ToBitPos(pos));
161 Pos pos_v = Direction::Shift(
ToBitPos(pos));
166 #endif // defined(__CUDA_ARCH__)
169 pos_v = Direction::Shift(pos_v);
171 value_type const value = bits_[pos_v.int_pos];
174 return static_cast<bool>(result);
185 inline friend std::ostream &
187 os <<
"Bits " <<
"storage size: " << field.bits_.size() <<
"\n";
189 std::bitset<BitFieldContainer<VT, Direction, IsConst>::kValueSize> bset(field.bits_[i]);
197 template <
typename VT,
bool IsConst = false>
211 template <
typename VT>
214 using Pos =
typename Container::Pos;
231 #endif // XGBOOST_COMMON_BITFIELD_H_
static constexpr value_type kOne
Definition: bitfield.h:64
XGBOOST_DEVICE pointer Data() const
Definition: bitfield.h:183
BitFieldContainer()=default
typename Container::value_type value_type
Definition: bitfield.h:201
value_type * pointer
Definition: bitfield.h:61
void Clear(value_type pos)
Definition: bitfield.h:160
static XGBOOST_DEVICE size_t ComputeStorageSize(size_t size)
Definition: bitfield.h:97
A non-owning type with auxiliary methods defined for manipulating bits.
Definition: bitfield.h:59
std::size_t index_type
Definition: span.h:416
common::Span< value_type > Bits()
Definition: bitfield.h:91
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:212
common::Span< value_type const > Bits() const
Definition: bitfield.h:92
BitFieldContainer & operator&=(BitFieldContainer const &rhs)
Definition: bitfield.h:129
typename Container::Pos Pos
Definition: bitfield.h:214
XGBOOST_DEVICE bool Check(Pos pos_v) const
Definition: bitfield.h:168
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:217
typename Container::Pos Pos
Definition: bitfield.h:200
Definition: bitfield.h:198
XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b)
Definition: common.h:84
static XGBOOST_DEVICE Pos Shift(Pos pos)
Definition: bitfield.h:203
Definition: bitfield.h:66
static constexpr value_type kValueSize
Definition: bitfield.h:63
#define SPAN_LT(lhs, rhs)
Definition: span.h:112
static XGBOOST_DEVICE Pos ToBitPos(value_type pos)
Definition: bitfield.h:76
void Set(value_type pos)
Definition: bitfield.h:154
typename Container::value_type value_type
Definition: bitfield.h:215
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:542
constexpr XGBOOST_DEVICE pointer data() const __span_noexcept
Definition: span.h:537
BitFieldContainer & operator|=(BitFieldContainer const &rhs)
Definition: bitfield.h:110
XGBOOST_DEVICE size_t Size() const
Definition: bitfield.h:181
XGBOOST_DEVICE BitFieldContainer(common::Span< value_type > bits)
Definition: bitfield.h:88
friend std::ostream & operator<<(std::ostream &os, BitFieldContainer< VT, Direction, IsConst > field)
Definition: bitfield.h:186
#define XGBOOST_DEVICE
Tag function as usable by device.
Definition: base.h:84
XGBOOST_DEVICE bool Check(value_type pos) const
Definition: bitfield.h:176
namespace of xgboost
Definition: base.h:110