diff --git a/src/tree/gpu_hist/evaluate_splits.cuh b/src/tree/gpu_hist/evaluate_splits.cuh index aebaa89b2e7b..e5105183c1a6 100644 --- a/src/tree/gpu_hist/evaluate_splits.cuh +++ b/src/tree/gpu_hist/evaluate_splits.cuh @@ -68,7 +68,7 @@ class GPUHistEvaluator { // storage for sorted index of feature histogram, used for sort based splits. dh::device_vector cat_sorted_idx_; // cached input for sorting the histogram, used for sort based splits. - using SortPair = thrust::tuple; + using SortPair = thrust::tuple; dh::device_vector sort_input_; // cache for feature index dh::device_vector feature_idx_; diff --git a/src/tree/gpu_hist/evaluator.cu b/src/tree/gpu_hist/evaluator.cu index 93d7215a2629..634734943aa0 100644 --- a/src/tree/gpu_hist/evaluator.cu +++ b/src/tree/gpu_hist/evaluator.cu @@ -89,7 +89,7 @@ common::Span GPUHistEvaluator::SortHistogram( input.gradient_histogram[j]); return thrust::make_tuple(i, lw); } - return thrust::make_tuple(i, 0.0); + return thrust::make_tuple(i, 0.0f); }); // Sort an array segmented according to // - nodes diff --git a/src/tree/split_evaluator.h b/src/tree/split_evaluator.h index a44522f88197..30263a6eb22d 100644 --- a/src/tree/split_evaluator.h +++ b/src/tree/split_evaluator.h @@ -66,21 +66,21 @@ class TreeEvaluator { template struct SplitEvaluator { - common::Span constraints; - common::Span lower; - common::Span upper; + const int* constraints; + const float* lower; + const float* upper; bool has_constraint; - XGBOOST_DEVICE double CalcSplitGain(const ParamT ¶m, bst_node_t nidx, + XGBOOST_DEVICE float CalcSplitGain(const ParamT ¶m, bst_node_t nidx, bst_feature_t fidx, tree::GradStats const& left, tree::GradStats const& right) const { int constraint = constraints[fidx]; - const double negative_infinity = -std::numeric_limits::infinity(); - double wleft = this->CalcWeight(nidx, param, left); - double wright = this->CalcWeight(nidx, param, right); + const float negative_infinity = -std::numeric_limits::infinity(); + float wleft = this->CalcWeight(nidx, param, left); + float wright = this->CalcWeight(nidx, param, right); - double gain = this->CalcGainGivenWeight(param, left, wleft) + + float gain = this->CalcGainGivenWeight(param, left, wleft) + this->CalcGainGivenWeight(param, right, wright); if (constraint == 0) { @@ -101,9 +101,9 @@ class TreeEvaluator { if (nodeid == kRootParentId) { return w; - } else if (w < lower(nodeid)) { + } else if (w < lower[nodeid]) { return lower[nodeid]; - } else if (w > upper(nodeid)) { + } else if (w > upper[nodeid]) { return upper[nodeid]; } else { return w; @@ -111,7 +111,7 @@ class TreeEvaluator { } template - XGBOOST_DEVICE double CalcWeightCat(ParamT const& param, GradientSumT const& stats) const { + XGBOOST_DEVICE float CalcWeightCat(ParamT const& param, GradientSumT const& stats) const { // FIXME(jiamingy): This is a temporary solution until we have categorical feature // specific regularization parameters. During sorting we should try to avoid any // regularization. @@ -141,15 +141,13 @@ class TreeEvaluator { /* Get a view to the evaluator that can be passed down to device. */ template auto GetEvaluator() const { if (device_ != GenericParameter::kCpuId) { - auto constraints = monotone_.ConstDeviceSpan(); - return SplitEvaluator{ - constraints, lower_bounds_.ConstDeviceSpan(), - upper_bounds_.ConstDeviceSpan(), has_constraint_}; + auto constraints = monotone_.ConstDevicePointer(); + return SplitEvaluator{constraints, lower_bounds_.ConstDevicePointer(), + upper_bounds_.ConstDevicePointer(), has_constraint_}; } else { - auto constraints = monotone_.ConstHostSpan(); - return SplitEvaluator{constraints, lower_bounds_.ConstHostSpan(), - upper_bounds_.ConstHostSpan(), - has_constraint_}; + auto constraints = monotone_.ConstHostPointer(); + return SplitEvaluator{constraints, lower_bounds_.ConstHostPointer(), + upper_bounds_.ConstHostPointer(), has_constraint_}; } }