From 9c469b3844855fd4139547fb5d8cf41eea29498c Mon Sep 17 00:00:00 2001 From: Jiaming Yuan Date: Tue, 6 Aug 2019 02:49:32 -0400 Subject: [PATCH] Move bitfield into common. (#4737) * Prepare for columnar format support. --- src/common/bitfield.cuh | 160 ++++++++++++++++++++ src/tree/constraints.cu | 1 + src/tree/constraints.cuh | 153 +------------------ tests/cpp/{tree => common}/test_bitfield.cu | 2 +- 4 files changed, 163 insertions(+), 153 deletions(-) create mode 100644 src/common/bitfield.cuh rename tests/cpp/{tree => common}/test_bitfield.cu (97%) diff --git a/src/common/bitfield.cuh b/src/common/bitfield.cuh new file mode 100644 index 000000000..8349e0198 --- /dev/null +++ b/src/common/bitfield.cuh @@ -0,0 +1,160 @@ +#ifndef XGBOOST_COMMON_BITFIELD_CUH_ +#define XGBOOST_COMMON_BITFIELD_CUH_ + +#include +#include +#include +#include +#include + +#include "span.h" + +namespace xgboost { + +__forceinline__ __device__ unsigned long long AtomicOr(unsigned long long* address, + unsigned long long val) { + unsigned long long int old = *address, assumed; // NOLINT + do { + assumed = old; + old = atomicCAS(address, assumed, val | assumed); + } while (assumed != old); + + return old; +} + +__forceinline__ __device__ unsigned long long AtomicAnd(unsigned long long* address, + unsigned long long val) { + unsigned long long int old = *address, assumed; // NOLINT + do { + assumed = old; + old = atomicCAS(address, assumed, val & assumed); + } while (assumed != old); + + return old; +} + +/*! + * \brief A non-owning type with auxiliary methods defined for manipulating bits. + */ +struct BitField { + using value_type = uint64_t; + + static value_type constexpr kValueSize = sizeof(value_type) * 8; + static value_type constexpr kOne = 1UL; // force uint64_t + static_assert(kValueSize == 64, "uint64_t should be of 64 bits."); + + struct Pos { + value_type int_pos {0}; + value_type bit_pos {0}; + }; + + common::Span bits_; + + public: + BitField() = default; + XGBOOST_DEVICE BitField(common::Span bits) : bits_{bits} {} + XGBOOST_DEVICE BitField(BitField const& other) : bits_{other.bits_} {} + + static size_t ComputeStorageSize(size_t size) { + auto pos = ToBitPos(size); + if (size < kValueSize) { + return 1; + } + + if (pos.bit_pos != 0) { + return pos.int_pos + 2; + } else { + return pos.int_pos + 1; + } + } + XGBOOST_DEVICE static Pos ToBitPos(value_type pos) { + Pos pos_v; + if (pos == 0) { + return pos_v; + } + pos_v.int_pos = pos / kValueSize; + pos_v.bit_pos = pos % kValueSize; + return pos_v; + } + + __device__ BitField& operator|=(BitField const& rhs) { + auto tid = blockIdx.x * blockDim.x + threadIdx.x; + size_t min_size = min(bits_.size(), rhs.bits_.size()); + if (tid < min_size) { + bits_[tid] |= rhs.bits_[tid]; + } + return *this; + } + __device__ BitField& operator&=(BitField const& rhs) { + size_t min_size = min(bits_.size(), rhs.bits_.size()); + auto tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid < min_size) { + bits_[tid] &= rhs.bits_[tid]; + } + return *this; + } + + XGBOOST_DEVICE size_t Size() const { return kValueSize * bits_.size(); } + + __device__ void Set(value_type pos) { + Pos pos_v = ToBitPos(pos); + value_type& value = bits_[pos_v.int_pos]; + value_type set_bit = kOne << (kValueSize - pos_v.bit_pos - kOne); + static_assert(sizeof(unsigned long long int) == sizeof(value_type), ""); + AtomicOr(reinterpret_cast(&value), set_bit); + } + __device__ void Clear(value_type pos) { + Pos pos_v = ToBitPos(pos); + value_type& value = bits_[pos_v.int_pos]; + value_type clear_bit = ~(kOne << (kValueSize - pos_v.bit_pos - kOne)); + static_assert(sizeof(unsigned long long int) == sizeof(value_type), ""); + AtomicAnd(reinterpret_cast(&value), clear_bit); + } + + XGBOOST_DEVICE bool Check(Pos pos_v) const { + value_type value = bits_[pos_v.int_pos]; + value_type const test_bit = kOne << (kValueSize - pos_v.bit_pos - kOne); + value_type result = test_bit & value; + return static_cast(result); + } + XGBOOST_DEVICE bool Check(value_type pos) const { + Pos pos_v = ToBitPos(pos); + return Check(pos_v); + } + + friend std::ostream& operator<<(std::ostream& os, BitField field) { + os << "Bits " << "storage size: " << field.bits_.size() << "\n"; + for (size_t i = 0; i < field.bits_.size(); ++i) { + std::bitset set(field.bits_[i]); + os << set << "\n"; + } + return os; + } +}; + +inline void PrintDeviceBits(std::string name, BitField field) { + std::cout << "Bits: " << name << std::endl; + std::vector h_field_bits(field.bits_.size()); + thrust::copy(thrust::device_ptr(field.bits_.data()), + thrust::device_ptr(field.bits_.data() + field.bits_.size()), + h_field_bits.data()); + BitField h_field; + h_field.bits_ = {h_field_bits.data(), h_field_bits.data() + h_field_bits.size()}; + std::cout << h_field; +} + +inline void PrintDeviceStorage(std::string name, common::Span list) { + std::cout << name << std::endl; + std::vector h_list(list.size()); + thrust::copy(thrust::device_ptr(list.data()), + thrust::device_ptr(list.data() + list.size()), + h_list.data()); + for (auto v : h_list) { + std::cout << v << ", "; + } + std::cout << std::endl; +} + +} + +#endif // XGBOOST_COMMON_BITFIELD_CUH_ \ No newline at end of file diff --git a/src/tree/constraints.cu b/src/tree/constraints.cu index 2000298bd..cba906bf2 100644 --- a/src/tree/constraints.cu +++ b/src/tree/constraints.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include "constraints.cuh" #include "param.h" diff --git a/src/tree/constraints.cuh b/src/tree/constraints.cuh index eebbea5da..e30530c70 100644 --- a/src/tree/constraints.cuh +++ b/src/tree/constraints.cuh @@ -5,168 +5,17 @@ #define XGBOOST_TREE_CONSTRAINTS_H_ #include -#include #include -#include #include -#include -#include -#include -#include #include "param.h" #include "../common/span.h" +#include "../common/bitfield.cuh" #include "../common/device_helpers.cuh" -#include - namespace xgboost { -__forceinline__ __device__ unsigned long long AtomicOr(unsigned long long* address, - unsigned long long val) { - unsigned long long int old = *address, assumed; // NOLINT - do { - assumed = old; - old = atomicCAS(address, assumed, val | assumed); - } while (assumed != old); - - return old; -} - -__forceinline__ __device__ unsigned long long AtomicAnd(unsigned long long* address, - unsigned long long val) { - unsigned long long int old = *address, assumed; // NOLINT - do { - assumed = old; - old = atomicCAS(address, assumed, val & assumed); - } while (assumed != old); - - return old; -} - -/*! - * \brief A non-owning type with auxiliary methods defined for manipulating bits. - */ -struct BitField { - using value_type = uint64_t; - - static value_type constexpr kValueSize = sizeof(value_type) * 8; - static value_type constexpr kOne = 1UL; // force uint64_t - static_assert(kValueSize == 64, "uint64_t should be of 64 bits."); - - struct Pos { - value_type int_pos {0}; - value_type bit_pos {0}; - }; - - common::Span bits_; - - public: - BitField() = default; - XGBOOST_DEVICE BitField(common::Span bits) : bits_{bits} {} - XGBOOST_DEVICE BitField(BitField const& other) : bits_{other.bits_} {} - - static size_t ComputeStorageSize(size_t size) { - auto pos = ToBitPos(size); - if (size < kValueSize) { - return 1; - } - - if (pos.bit_pos != 0) { - return pos.int_pos + 2; - } else { - return pos.int_pos + 1; - } - } - XGBOOST_DEVICE static Pos ToBitPos(value_type pos) { - Pos pos_v; - if (pos == 0) { - return pos_v; - } - pos_v.int_pos = pos / kValueSize; - pos_v.bit_pos = pos % kValueSize; - return pos_v; - } - - __device__ BitField& operator|=(BitField const& rhs) { - auto tid = blockIdx.x * blockDim.x + threadIdx.x; - size_t min_size = min(bits_.size(), rhs.bits_.size()); - if (tid < min_size) { - bits_[tid] |= rhs.bits_[tid]; - } - return *this; - } - __device__ BitField& operator&=(BitField const& rhs) { - size_t min_size = min(bits_.size(), rhs.bits_.size()); - auto tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid < min_size) { - bits_[tid] &= rhs.bits_[tid]; - } - return *this; - } - - XGBOOST_DEVICE size_t Size() const { return kValueSize * bits_.size(); } - - __device__ void Set(value_type pos) { - Pos pos_v = ToBitPos(pos); - value_type& value = bits_[pos_v.int_pos]; - value_type set_bit = kOne << (kValueSize - pos_v.bit_pos - kOne); - static_assert(sizeof(unsigned long long int) == sizeof(value_type), ""); - AtomicOr(reinterpret_cast(&value), set_bit); - } - __device__ void Clear(value_type pos) { - Pos pos_v = ToBitPos(pos); - value_type& value = bits_[pos_v.int_pos]; - value_type clear_bit = ~(kOne << (kValueSize - pos_v.bit_pos - kOne)); - static_assert(sizeof(unsigned long long int) == sizeof(value_type), ""); - AtomicAnd(reinterpret_cast(&value), clear_bit); - } - - XGBOOST_DEVICE bool Check(Pos pos_v) const { - value_type value = bits_[pos_v.int_pos]; - value_type const test_bit = kOne << (kValueSize - pos_v.bit_pos - kOne); - value_type result = test_bit & value; - return static_cast(result); - } - XGBOOST_DEVICE bool Check(value_type pos) const { - Pos pos_v = ToBitPos(pos); - return Check(pos_v); - } - - friend std::ostream& operator<<(std::ostream& os, BitField field) { - os << "Bits " << "storage size: " << field.bits_.size() << "\n"; - for (size_t i = 0; i < field.bits_.size(); ++i) { - std::bitset set(field.bits_[i]); - os << set << "\n"; - } - return os; - } -}; - -inline void PrintDeviceBits(std::string name, BitField field) { - std::cout << "Bits: " << name << std::endl; - std::vector h_field_bits(field.bits_.size()); - thrust::copy(thrust::device_ptr(field.bits_.data()), - thrust::device_ptr(field.bits_.data() + field.bits_.size()), - h_field_bits.data()); - BitField h_field; - h_field.bits_ = {h_field_bits.data(), h_field_bits.data() + h_field_bits.size()}; - std::cout << h_field; -} - -inline void PrintDeviceStorage(std::string name, common::Span list) { - std::cout << name << std::endl; - std::vector h_list(list.size()); - thrust::copy(thrust::device_ptr(list.data()), - thrust::device_ptr(list.data() + list.size()), - h_list.data()); - for (auto v : h_list) { - std::cout << v << ", "; - } - std::cout << std::endl; -} - // Feature interaction constraints built for GPU Hist updater. struct FeatureInteractionConstraint { protected: diff --git a/tests/cpp/tree/test_bitfield.cu b/tests/cpp/common/test_bitfield.cu similarity index 97% rename from tests/cpp/tree/test_bitfield.cu rename to tests/cpp/common/test_bitfield.cu index aa5d36d49..ad78da2e2 100644 --- a/tests/cpp/tree/test_bitfield.cu +++ b/tests/cpp/common/test_bitfield.cu @@ -5,7 +5,7 @@ #include #include #include -#include "../../../src/tree/constraints.cuh" +#include "../../../src/common/bitfield.cuh" #include "../../../src/common/device_helpers.cuh" namespace xgboost {