Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Use quantised gradients in gpu_hist histograms (#8246)

Browse Source
This commit is contained in:
Rory Mitchell
2022-09-26 17:35:35 +02:00
committed by GitHub
parent 4056974e37
commit 8f77677193
14 changed files with 394 additions and 336 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
src/tree/gpu_hist/evaluate_splits.cu
View File

@@ -58,7 +58,8 @@ class EvaluateSplitAgent {
const uint32_t gidx_begin; // beginning bin
const uint32_t gidx_end; // end bin for i^th feature
const dh::LDGIterator<float> feature_values;
const GradientPairPrecise *node_histogram;
const GradientPairInt64 *node_histogram;
const GradientQuantizer &rounding;
const GradientPairPrecise parent_sum;
const GradientPairPrecise missing;
const GPUTrainingParam &param;
@@ -79,6 +80,7 @@ class EvaluateSplitAgent {
gidx_end(__ldg(shared_inputs.feature_segments.data() + fidx + 1)),
feature_values(shared_inputs.feature_values.data()),
node_histogram(inputs.gradient_histogram.data()),
rounding(shared_inputs.rounding),
parent_sum(dh::LDGIterator<GradientPairPrecise>(&inputs.parent_sum)[0]),
param(shared_inputs.param),
evaluator(evaluator),
@@ -98,11 +100,12 @@ class EvaluateSplitAgent {
}
// Load using efficient 128 vector load instruction
__device__ __forceinline__ GradientPairPrecise LoadGpair(const GradientPairPrecise *ptr) {
static_assert(sizeof(GradientPairPrecise) == sizeof(float4),
"Vector type size does not match gradient pair size.");
__device__ __forceinline__ GradientPairPrecise LoadGpair(const GradientPairInt64 *ptr) {
float4 tmp = *reinterpret_cast<const float4 *>(ptr);
return *reinterpret_cast<const GradientPairPrecise *>(&tmp);
auto gpair_int = *reinterpret_cast<const GradientPairInt64 *>(&tmp);
static_assert(sizeof(decltype(gpair_int)) == sizeof(float4),
"Vector type size does not match gradient pair size.");
return rounding.ToFloatingPoint(gpair_int);
}
__device__ __forceinline__ void Numerical(DeviceSplitCandidate *__restrict__ best_split) {

4
src/tree/gpu_hist/evaluate_splits.cuh
View File

@@ -10,6 +10,7 @@
#include "../split_evaluator.h"
#include "../updater_gpu_common.cuh"
#include "expand_entry.cuh"
#include "histogram.cuh"
namespace xgboost {
namespace common {
@@ -24,12 +25,13 @@ struct EvaluateSplitInputs {
int depth;
GradientPairPrecise parent_sum;
common::Span<const bst_feature_t> feature_set;
common::Span<const GradientPairPrecise> gradient_histogram;
common::Span<const GradientPairInt64> gradient_histogram;
};
// Inputs necessary for all nodes
struct EvaluateSplitSharedInputs {
GPUTrainingParam param;
GradientQuantizer rounding;
common::Span<FeatureType const> feature_types;
common::Span<const uint32_t> feature_segments;
common::Span<const float> feature_values;

5
src/tree/gpu_hist/evaluator.cu
View File

@@ -83,8 +83,9 @@ common::Span<bst_feature_t const> GPUHistEvaluator::SortHistogram(
auto j = i % total_bins;
auto fidx = d_feature_idx[j];
if (common::IsCat(shared_inputs.feature_types, fidx)) {
auto lw = evaluator.CalcWeightCat(shared_inputs.param,
input.gradient_histogram[j]);
auto grad =
shared_inputs.rounding.ToFloatingPoint(input.gradient_histogram[j]);
auto lw = evaluator.CalcWeightCat(shared_inputs.param, grad);
return thrust::make_tuple(i, lw);
}
return thrust::make_tuple(i, 0.0f);

129
src/tree/gpu_hist/histogram.cu
View File

@@ -72,30 +72,35 @@ struct Clip : public thrust::unary_function<GradientPair, Pair> {
}
};
template <typename GradientSumT>
HistRounding<GradientSumT> CreateRoundingFactor(common::Span<GradientPair const> gpair) {
GradientQuantizer::GradientQuantizer(common::Span<GradientPair const> gpair) {
using GradientSumT = GradientPairPrecise;
using T = typename GradientSumT::ValueT;
dh::XGBCachingDeviceAllocator<char> alloc;
thrust::device_ptr<GradientPair const> gpair_beg{gpair.data()};
thrust::device_ptr<GradientPair const> gpair_end{gpair.data() + gpair.size()};
auto beg = thrust::make_transform_iterator(gpair_beg, Clip());
auto end = thrust::make_transform_iterator(gpair_end, Clip());
Pair p = dh::Reduce(thrust::cuda::par(alloc), beg, end, Pair{}, thrust::plus<Pair>{});
Pair p =
dh::Reduce(thrust::cuda::par(alloc), beg, beg + gpair.size(), Pair{}, thrust::plus<Pair>{});
// Treat pair as array of 4 primitive types to allreduce
using ReduceT = typename decltype(p.first)::ValueT;
static_assert(sizeof(Pair) == sizeof(ReduceT) * 4, "Expected to reduce four elements.");
rabit::Allreduce<rabit::op::Sum, ReduceT>(reinterpret_cast<ReduceT*>(&p), 4);
GradientPair positive_sum{p.first}, negative_sum{p.second};
auto histogram_rounding =
GradientSumT{CreateRoundingFactor<T>(std::max(positive_sum.GetGrad(), negative_sum.GetGrad()),
gpair.size()),
CreateRoundingFactor<T>(std::max(positive_sum.GetHess(), negative_sum.GetHess()),
gpair.size())};
std::size_t total_rows = gpair.size();
rabit::Allreduce<rabit::op::Sum>(&total_rows, 1);
using IntT = typename HistRounding<GradientSumT>::SharedSumT::ValueT;
auto histogram_rounding = GradientSumT{
CreateRoundingFactor<T>(std::max(positive_sum.GetGrad(), negative_sum.GetGrad()), total_rows),
CreateRoundingFactor<T>(std::max(positive_sum.GetHess(), negative_sum.GetHess()),
total_rows)};
using IntT = typename GradientPairInt64::ValueT;
/**
* Factor for converting gradients from fixed-point to floating-point.
*/
GradientSumT to_floating_point =
to_floating_point_ =
histogram_rounding /
T(IntT(1) << (sizeof(typename GradientSumT::ValueT) * 8 - 2)); // keep 1 for sign bit
/**
@@ -107,35 +112,55 @@ HistRounding<GradientSumT> CreateRoundingFactor(common::Span<GradientPair const>
* rounding is calcuated as exp(m), see the rounding factor calcuation for
* details.
*/
GradientSumT to_fixed_point =
GradientSumT(T(1) / to_floating_point.GetGrad(), T(1) / to_floating_point.GetHess());
return {histogram_rounding, to_fixed_point, to_floating_point};
to_fixed_point_ =
GradientSumT(T(1) / to_floating_point_.GetGrad(), T(1) / to_floating_point_.GetHess());
}
template HistRounding<GradientPairPrecise> CreateRoundingFactor(
common::Span<GradientPair const> gpair);
template HistRounding<GradientPair> CreateRoundingFactor(common::Span<GradientPair const> gpair);
template <typename GradientSumT, int kBlockThreads, int kItemsPerThread,
XGBOOST_DEV_INLINE void
AtomicAddGpairShared(xgboost::GradientPairInt64 *dest,
xgboost::GradientPairInt64 const &gpair) {
auto dst_ptr = reinterpret_cast<int64_t *>(dest);
auto g = gpair.GetQuantisedGrad();
auto h = gpair.GetQuantisedHess();
AtomicAdd64As32(dst_ptr, g);
AtomicAdd64As32(dst_ptr + 1, h);
}
// Global 64 bit integer atomics at the time of writing do not benefit from being separated into two
// 32 bit atomics
XGBOOST_DEV_INLINE void AtomicAddGpairGlobal(xgboost::GradientPairInt64* dest,
xgboost::GradientPairInt64 const& gpair) {
auto dst_ptr = reinterpret_cast<uint64_t*>(dest);
auto g = gpair.GetQuantisedGrad();
auto h = gpair.GetQuantisedHess();
atomicAdd(dst_ptr,
*reinterpret_cast<uint64_t*>(&g));
atomicAdd(dst_ptr + 1,
*reinterpret_cast<uint64_t*>(&h));
}
template <int kBlockThreads, int kItemsPerThread,
int kItemsPerTile = kBlockThreads* kItemsPerThread>
class HistogramAgent {
using SharedSumT = typename HistRounding<GradientSumT>::SharedSumT;
SharedSumT* smem_arr_;
GradientSumT* d_node_hist_;
GradientPairInt64* smem_arr_;
GradientPairInt64* d_node_hist_;
dh::LDGIterator<const RowPartitioner::RowIndexT> d_ridx_;
const GradientPair* d_gpair_;
const FeatureGroup group_;
const EllpackDeviceAccessor& matrix_;
const int feature_stride_;
const std::size_t n_elements_;
const HistRounding<GradientSumT>& rounding_;
const GradientQuantizer& rounding_;
public:
__device__ HistogramAgent(SharedSumT* smem_arr, GradientSumT* __restrict__ d_node_hist,
const FeatureGroup& group, const EllpackDeviceAccessor& matrix,
__device__ HistogramAgent(GradientPairInt64* smem_arr,
GradientPairInt64* __restrict__ d_node_hist, const FeatureGroup& group,
const EllpackDeviceAccessor& matrix,
common::Span<const RowPartitioner::RowIndexT> d_ridx,
const HistRounding<GradientSumT>& rounding, const GradientPair* d_gpair)
const GradientQuantizer& rounding, const GradientPair* d_gpair)
: smem_arr_(smem_arr),
d_node_hist_(d_node_hist),
d_ridx_(d_ridx.data()),
@@ -155,7 +180,7 @@ class HistogramAgent {
group_.start_bin;
if (matrix_.is_dense || gidx != matrix_.NumBins()) {
auto adjusted = rounding_.ToFixedPoint(d_gpair_[ridx]);
dh::AtomicAddGpair(smem_arr_ + gidx, adjusted);
AtomicAddGpairShared(smem_arr_ + gidx, adjusted);
}
}
}
@@ -185,12 +210,12 @@ class HistogramAgent {
for (int i = 0; i < kItemsPerThread; i++) {
if ((matrix_.is_dense || gidx[i] != matrix_.NumBins())) {
auto adjusted = rounding_.ToFixedPoint(gpair[i]);
dh::AtomicAddGpair(smem_arr_ + gidx[i] - group_.start_bin, adjusted);
AtomicAddGpairShared(smem_arr_ + gidx[i] - group_.start_bin, adjusted);
}
}
}
__device__ void BuildHistogramWithShared() {
dh::BlockFill(smem_arr_, group_.num_bins, SharedSumT());
dh::BlockFill(smem_arr_, group_.num_bins, GradientPairInt64());
__syncthreads();
std::size_t offset = blockIdx.x * kItemsPerTile;
@@ -203,8 +228,7 @@ class HistogramAgent {
// Write shared memory back to global memory
__syncthreads();
for (auto i : dh::BlockStrideRange(0, group_.num_bins)) {
auto truncated = rounding_.ToFloatingPoint(smem_arr_[i]);
dh::AtomicAddGpair(d_node_hist_ + group_.start_bin + i, truncated);
AtomicAddGpairGlobal(d_node_hist_ + group_.start_bin + i, smem_arr_[i]);
}
}
@@ -215,36 +239,26 @@ class HistogramAgent {
matrix_
.gidx_iter[ridx * matrix_.row_stride + group_.start_feature + idx % feature_stride_];
if (matrix_.is_dense || gidx != matrix_.NumBins()) {
// If we are not using shared memory, accumulate the values directly into
// global memory
GradientSumT truncated{
TruncateWithRoundingFactor<GradientSumT::ValueT>(rounding_.rounding.GetGrad(),
d_gpair_[ridx].GetGrad()),
TruncateWithRoundingFactor<GradientSumT::ValueT>(rounding_.rounding.GetHess(),
d_gpair_[ridx].GetHess()),
};
dh::AtomicAddGpair(d_node_hist_ + gidx, truncated);
auto adjusted = rounding_.ToFixedPoint(d_gpair_[ridx]);
AtomicAddGpairGlobal(d_node_hist_ + gidx, adjusted);
}
}
}
};
template <typename GradientSumT, bool use_shared_memory_histograms, int kBlockThreads,
template <bool use_shared_memory_histograms, int kBlockThreads,
int kItemsPerThread>
__global__ void __launch_bounds__(kBlockThreads)
SharedMemHistKernel(const EllpackDeviceAccessor matrix,
const FeatureGroupsAccessor feature_groups,
common::Span<const RowPartitioner::RowIndexT> d_ridx,
GradientSumT* __restrict__ d_node_hist,
GradientPairInt64* __restrict__ d_node_hist,
const GradientPair* __restrict__ d_gpair,
HistRounding<GradientSumT> const rounding) {
using SharedSumT = typename HistRounding<GradientSumT>::SharedSumT;
using T = typename GradientSumT::ValueT;
GradientQuantizer const rounding) {
extern __shared__ char smem[];
const FeatureGroup group = feature_groups[blockIdx.y];
SharedSumT* smem_arr = reinterpret_cast<SharedSumT*>(smem);
auto agent = HistogramAgent<GradientSumT, kBlockThreads, kItemsPerThread>(
auto smem_arr = reinterpret_cast<GradientPairInt64*>(smem);
auto agent = HistogramAgent<kBlockThreads, kItemsPerThread>(
smem_arr, d_node_hist, group, matrix, d_ridx, rounding, d_gpair);
if (use_shared_memory_histograms) {
agent.BuildHistogramWithShared();
@@ -253,13 +267,12 @@ __global__ void __launch_bounds__(kBlockThreads)
}
}
template <typename GradientSumT>
void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
FeatureGroupsAccessor const& feature_groups,
common::Span<GradientPair const> gpair,
common::Span<const uint32_t> d_ridx,
common::Span<GradientSumT> histogram,
HistRounding<GradientSumT> rounding, bool force_global_memory) {
common::Span<GradientPairInt64> histogram,
GradientQuantizer rounding, bool force_global_memory) {
// decide whether to use shared memory
int device = 0;
dh::safe_cuda(cudaGetDevice(&device));
@@ -267,7 +280,7 @@ void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
size_t max_shared_memory = dh::MaxSharedMemoryOptin(device);
size_t smem_size =
sizeof(typename HistRounding<GradientSumT>::SharedSumT) * feature_groups.max_group_bins;
sizeof(GradientPairInt64) * feature_groups.max_group_bins;
bool shared = !force_global_memory && smem_size <= max_shared_memory;
smem_size = shared ? smem_size : 0;
@@ -311,19 +324,13 @@ void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
};
if (shared) {
runit(SharedMemHistKernel<GradientSumT, true, kBlockThreads, kItemsPerThread>);
runit(SharedMemHistKernel<true, kBlockThreads, kItemsPerThread>);
} else {
runit(SharedMemHistKernel<GradientSumT, false, kBlockThreads, kItemsPerThread>);
runit(SharedMemHistKernel<false, kBlockThreads, kItemsPerThread>);
}
dh::safe_cuda(cudaGetLastError());
}
template void BuildGradientHistogram<GradientPairPrecise>(
EllpackDeviceAccessor const& matrix, FeatureGroupsAccessor const& feature_groups,
common::Span<GradientPair const> gpair, common::Span<const uint32_t> ridx,
common::Span<GradientPairPrecise> histogram, HistRounding<GradientPairPrecise> rounding,
bool force_global_memory);
} // namespace tree
} // namespace xgboost

69
src/tree/gpu_hist/histogram.cuh
View File

@@ -12,56 +12,51 @@
namespace xgboost {
namespace tree {
template <typename T, typename U>
XGBOOST_DEV_INLINE T TruncateWithRoundingFactor(T const rounding_factor, U const x) {
static_assert(sizeof(T) >= sizeof(U), "Rounding must have higher or equal precision.");
return (rounding_factor + static_cast<T>(x)) - rounding_factor;
/**
* \brief An atomicAdd designed for gradient pair with better performance. For general
* int64_t atomicAdd, one can simply cast it to unsigned long long. Exposed for testing.
*/
XGBOOST_DEV_INLINE void AtomicAdd64As32(int64_t* dst, int64_t src) {
uint32_t* y_low = reinterpret_cast<uint32_t*>(dst);
uint32_t* y_high = y_low + 1;
auto cast_src = reinterpret_cast<uint64_t *>(&src);
uint32_t const x_low = static_cast<uint32_t>(src);
uint32_t const x_high = (*cast_src) >> 32;
auto const old = atomicAdd(y_low, x_low);
uint32_t const carry = old > (std::numeric_limits<uint32_t>::max() - x_low) ? 1 : 0;
uint32_t const sig = x_high + carry;
atomicAdd(y_high, sig);
}
/**
* Truncation factor for gradient, see comments in `CreateRoundingFactor()` for details.
*/
template <typename GradientSumT>
struct HistRounding {
/* Factor to truncate the gradient before building histogram for deterministic result. */
GradientSumT rounding;
class GradientQuantizer {
private:
/* Convert gradient to fixed point representation. */
GradientSumT to_fixed_point;
GradientPairPrecise to_fixed_point_;
/* Convert fixed point representation back to floating point. */
GradientSumT to_floating_point;
/* Type used in shared memory. */
using SharedSumT = std::conditional_t<
std::is_same<typename GradientSumT::ValueT, float>::value,
GradientPairInt32, GradientPairInt64>;
using T = typename GradientSumT::ValueT;
XGBOOST_DEV_INLINE SharedSumT ToFixedPoint(GradientPair const& gpair) const {
auto adjusted = SharedSumT(T(gpair.GetGrad() * to_fixed_point.GetGrad()),
T(gpair.GetHess() * to_fixed_point.GetHess()));
GradientPairPrecise to_floating_point_;
public:
explicit GradientQuantizer(common::Span<GradientPair const> gpair);
XGBOOST_DEVICE GradientPairInt64 ToFixedPoint(GradientPair const& gpair) const {
auto adjusted = GradientPairInt64(gpair.GetGrad() * to_fixed_point_.GetGrad(),
gpair.GetHess() * to_fixed_point_.GetHess());
return adjusted;
}
XGBOOST_DEV_INLINE GradientSumT ToFloatingPoint(SharedSumT const &gpair) const {
auto g = gpair.GetGrad() * to_floating_point.GetGrad();
auto h = gpair.GetHess() * to_floating_point.GetHess();
GradientSumT truncated{
TruncateWithRoundingFactor<T>(rounding.GetGrad(), g),
TruncateWithRoundingFactor<T>(rounding.GetHess(), h),
};
return truncated;
XGBOOST_DEVICE GradientPairPrecise ToFloatingPoint(const GradientPairInt64&gpair) const {
auto g = gpair.GetQuantisedGrad() * to_floating_point_.GetGrad();
auto h = gpair.GetQuantisedHess() * to_floating_point_.GetHess();
return {g,h};
}
};
template <typename GradientSumT>
HistRounding<GradientSumT> CreateRoundingFactor(common::Span<GradientPair const> gpair);
template <typename GradientSumT>
void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
FeatureGroupsAccessor const& feature_groups,
common::Span<GradientPair const> gpair,
common::Span<const uint32_t> ridx,
common::Span<GradientSumT> histogram,
HistRounding<GradientSumT> rounding,
common::Span<GradientPairInt64> histogram,
GradientQuantizer rounding,
bool force_global_memory = false);
} // namespace tree
} // namespace xgboost

28
src/tree/updater_gpu_hist.cu
View File

@@ -72,9 +72,10 @@ DMLC_REGISTER_PARAMETER(GPUHistMakerTrainParam);
* \author Rory
* \date 28/07/2018
*/
template <typename GradientSumT, size_t kStopGrowingSize = 1 << 28>
template <size_t kStopGrowingSize = 1 << 28>
class DeviceHistogramStorage {
private:
using GradientSumT = GradientPairInt64;
/*! \brief Map nidx to starting index of its histogram. */
std::map<int, size_t> nidx_map_;
// Large buffer of zeroed memory, caches histograms
@@ -180,7 +181,7 @@ struct GPUHistMakerDevice {
BatchParam batch_param;
std::unique_ptr<RowPartitioner> row_partitioner;
DeviceHistogramStorage<GradientSumT> hist{};
DeviceHistogramStorage<> hist{};
dh::device_vector<GradientPair> d_gpair; // storage for gpair;
common::Span<GradientPair> gpair;
@@ -193,7 +194,7 @@ struct GPUHistMakerDevice {
TrainParam param;
HistRounding<GradientSumT> histogram_rounding;
std::unique_ptr<GradientQuantizer> histogram_rounding;
dh::PinnedMemory pinned;
dh::PinnedMemory pinned2;
@@ -265,7 +266,7 @@ struct GPUHistMakerDevice {
page = sample.page;
gpair = sample.gpair;
histogram_rounding = CreateRoundingFactor<GradientSumT>(this->gpair);
histogram_rounding.reset(new GradientQuantizer(this->gpair));
row_partitioner.reset(); // Release the device memory first before reallocating
row_partitioner.reset(new RowPartitioner(ctx_->gpu_id, sample.sample_rows));
@@ -282,7 +283,11 @@ struct GPUHistMakerDevice {
auto matrix = page->GetDeviceAccessor(ctx_->gpu_id);
EvaluateSplitInputs inputs{nidx, 0, root_sum, feature_set, hist.GetNodeHistogram(nidx)};
EvaluateSplitSharedInputs shared_inputs{
gpu_param, feature_types, matrix.feature_segments, matrix.gidx_fvalue_map,
gpu_param,
*histogram_rounding,
feature_types,
matrix.feature_segments,
matrix.gidx_fvalue_map,
matrix.min_fvalue,
};
auto split = this->evaluator_.EvaluateSingleSplit(inputs, shared_inputs);
@@ -298,7 +303,7 @@ struct GPUHistMakerDevice {
auto h_node_inputs = pinned2.GetSpan<EvaluateSplitInputs>(2 * candidates.size());
auto matrix = page->GetDeviceAccessor(ctx_->gpu_id);
EvaluateSplitSharedInputs shared_inputs{
GPUTrainingParam{param}, feature_types, matrix.feature_segments,
GPUTrainingParam{param}, *histogram_rounding, feature_types, matrix.feature_segments,
matrix.gidx_fvalue_map, matrix.min_fvalue,
};
dh::TemporaryArray<GPUExpandEntry> entries(2 * candidates.size());
@@ -344,7 +349,7 @@ struct GPUHistMakerDevice {
auto d_ridx = row_partitioner->GetRows(nidx);
BuildGradientHistogram(page->GetDeviceAccessor(ctx_->gpu_id),
feature_groups->DeviceAccessor(ctx_->gpu_id), gpair,
d_ridx, d_node_hist, histogram_rounding);
d_ridx, d_node_hist, *histogram_rounding);
}
// Attempt to do subtraction trick
@@ -526,11 +531,10 @@ struct GPUHistMakerDevice {
void AllReduceHist(int nidx, dh::AllReducer* reducer, int num_histograms) {
monitor.Start("AllReduce");
auto d_node_hist = hist.GetNodeHistogram(nidx).data();
reducer->AllReduceSum(reinterpret_cast<typename GradientSumT::ValueT*>(d_node_hist),
reinterpret_cast<typename GradientSumT::ValueT*>(d_node_hist),
page->Cuts().TotalBins() *
(sizeof(GradientSumT) / sizeof(typename GradientSumT::ValueT)) *
num_histograms);
using ReduceT = typename std::remove_pointer<decltype(d_node_hist)>::type::ValueT;
reducer->AllReduceSum(reinterpret_cast<ReduceT*>(d_node_hist),
reinterpret_cast<ReduceT*>(d_node_hist),
page->Cuts().TotalBins() * 2 * num_histograms);
monitor.Stop("AllReduce");
}