Skip to content

Commit

Permalink
Add debug logging
Browse files Browse the repository at this point in the history
  • Loading branch information
@hcho3
hcho3 committed Sep 7, 2021
1 parent ac530a5 commit b32a634
Show file tree
Hide file tree
Showing 5 changed files with 162 additions and 7 deletions.
97 changes: 93 additions & 4 deletions src/tree/gpu_hist/evaluate_splits.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -101,6 +101,74 @@ EvaluateSplitsFindOptimalSplitsViaScan(
auto scan_input_iter =
dh::MakeTransformIterator<thrust::tuple<ScanElem<GradientSumT>, ScanElem<GradientSumT>>>(
zip_loc_iter, ScanValueOp<GradientSumT>{left, right, evaluator});
{
auto write_computed_result = [](std::ostream& os, const ScanComputedElem<GradientSumT>& m) {
std::string best_direction_str =
(m.best_direction == DefaultDirection::kLeftDir) ? "left" : "right";
os << "(left_sum: " << m.left_sum << ", right_sum: " << m.right_sum
<< ", parent_sum: " << m.parent_sum << ", best_loss_chg: " << m.best_loss_chg
<< ", best_findex: " << m.best_findex << ", best_fvalue: " << m.best_fvalue
<< ", best_direction: " << best_direction_str << ")";
};
auto write_scan_elem = [&](std::ostream& os, const ScanElem<GradientSumT>& m) {
std::string indicator_str =
(m.indicator == ChildNodeIndicator::kLeftChild) ? "kLeftChild" : "kRightChild";
os << "(head_flag: " << (m.head_flag ? "true" : "false") << ", indicator: " << indicator_str
<< ", hist_idx: " << m.hist_idx << ", findex: " << m.findex << ", gpair: "<< m.gpair
<< ", fvalue: " << m.fvalue << ", is_cat: " << (m.is_cat ? "true" : "false")
<< ", computed_result: ";
write_computed_result(os, m.computed_result);
os << ")";
};
{
using TupT = thrust::tuple<ScanElem<GradientSumT>, ScanElem<GradientSumT>>;
thrust::device_vector<TupT> d_vec(size);
thrust::host_vector<TupT> vec(size);
thrust::copy(thrust::device, scan_input_iter, scan_input_iter + size, d_vec.begin());
thrust::copy(d_vec.begin(), d_vec.end(), vec.begin());
std::ostringstream oss;
for (const auto& e: vec) {
auto fw = thrust::get<0>(e);
auto bw = thrust::get<1>(e);
oss << "forward: ";
write_scan_elem(oss, fw);
oss << std::endl;
oss << "backward: ";
write_scan_elem(oss, bw);
oss << std::endl;
}
LOG(CONSOLE) << oss.str();
}

{
using TupT = thrust::tuple<ScanElem<GradientSumT>, ScanElem<GradientSumT>>;
thrust::device_vector<TupT> d_vec(size);

auto scan_op = ScanOp<GradientSumT>{left, right, evaluator};
std::size_t n_temp_bytes = 0;
cub::DeviceScan::InclusiveScan(nullptr, n_temp_bytes, scan_input_iter, d_vec.begin(),
scan_op, size);
dh::TemporaryArray<int8_t> temp(n_temp_bytes);
cub::DeviceScan::InclusiveScan(temp.data().get(), n_temp_bytes, scan_input_iter,
d_vec.begin(),
scan_op, size);

thrust::host_vector<TupT> vec(size);
thrust::copy(d_vec.begin(), d_vec.end(), vec.begin());
std::ostringstream oss;
for (const auto& e: vec) {
auto fw = thrust::get<0>(e);
auto bw = thrust::get<1>(e);
oss << "forward: ";
write_scan_elem(oss, fw);
oss << std::endl;
oss << "backward: ";
write_scan_elem(oss, bw);
oss << std::endl;
}
LOG(CONSOLE) << oss.str();
}
}

dh::device_vector<ScanComputedElem<GradientSumT>> out_scan(l_n_features + r_n_features);
auto scan_out_iter = thrust::make_transform_output_iterator(
Expand All @@ -114,6 +182,24 @@ EvaluateSplitsFindOptimalSplitsViaScan(
dh::TemporaryArray<int8_t> temp(n_temp_bytes);
cub::DeviceScan::InclusiveScan(temp.data().get(), n_temp_bytes, scan_input_iter, scan_out_iter,
scan_op, size);
{
auto write_computed_result = [](std::ostream& os, const ScanComputedElem<GradientSumT>& m) {
std::string best_direction_str =
(m.best_direction == DefaultDirection::kLeftDir) ? "left" : "right";
os << "(left_sum: " << m.left_sum << ", right_sum: " << m.right_sum
<< ", parent_sum: " << m.parent_sum << ", best_loss_chg: " << m.best_loss_chg
<< ", best_findex: " << m.best_findex << ", best_fvalue: " << m.best_fvalue
<< ", best_direction: " << best_direction_str << ")";
};
thrust::host_vector<ScanComputedElem<GradientSumT>> h_out_scan(l_n_features + r_n_features);
thrust::copy(out_scan.begin(), out_scan.end(), h_out_scan.begin());
std::ostringstream oss;
for (const auto& e : h_out_scan) {
write_computed_result(oss, e);
oss << std::endl;
}
LOG(CONSOLE) << oss.str();
}
return out_scan;
}

Expand All @@ -136,6 +222,7 @@ ScanValueOp<GradientSumT>::MapEvaluateSplitsHistEntryToScanElem(
* For the element at the beginning of each segment, compute gradient sums and loss_chg
* ahead of time. These will be later used by the inclusive scan operator.
**/
ret.head_flag = true;
if (ret.is_cat) {
ret.computed_result.left_sum = split_input.parent_sum - ret.gpair;
ret.computed_result.right_sum = ret.gpair;
Expand All @@ -159,6 +246,8 @@ ScanValueOp<GradientSumT>::MapEvaluateSplitsHistEntryToScanElem(
ret.computed_result.best_fvalue = ret.fvalue;
ret.computed_result.best_direction =
(forward ? DefaultDirection::kRightDir : DefaultDirection::kLeftDir);
} else {
ret.head_flag = false;
}

return ret;
Expand All @@ -184,13 +273,13 @@ template <typename GradientSumT>
template <bool forward>
__noinline__ __device__ ScanElem<GradientSumT>
ScanOp<GradientSumT>::DoIt(ScanElem<GradientSumT> lhs, ScanElem<GradientSumT> rhs) {
ScanElem<GradientSumT> ret;
ret = rhs;
ret.computed_result = {};
if (lhs.findex != rhs.findex || lhs.indicator != rhs.indicator) {
if (rhs.head_flag) {
// Segmented Scan
return rhs;
}
ScanElem<GradientSumT> ret;
ret = rhs;
ret.head_flag = (lhs.head_flag || rhs.head_flag);
if (((lhs.indicator == ChildNodeIndicator::kLeftChild) &&
(left.feature_set.size() != left.feature_segments.size()) &&
!thrust::binary_search(thrust::seq, left.feature_set.begin(), left.feature_set.end(),
Expand Down
3 changes: 2 additions & 1 deletion src/tree/gpu_hist/evaluate_splits.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,8 @@ struct ScanComputedElem {
template <typename GradientSumT>
struct ScanElem {
ChildNodeIndicator indicator{ChildNodeIndicator::kLeftChild};
uint64_t hist_idx;
uint64_t hist_idx{0};
bool head_flag{false};
GradientSumT gpair{0.0, 0.0};
int32_t findex{-1};
float fvalue{std::numeric_limits<float>::quiet_NaN()};
Expand Down
3 changes: 2 additions & 1 deletion tests/cpp/histogram_helpers.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,8 @@ class HistogramCutsWrapper : public common::HistogramCuts {
} // anonymous namespace

inline std::unique_ptr<EllpackPageImpl> BuildEllpackPage(
int n_rows, int n_cols, bst_float sparsity= 0) {
int n_rows, int n_cols, bst_float sparsity = 0) {
CHECK_EQ(n_cols, 8) << "n_cols must be equal to 8";
auto dmat = RandomDataGenerator(n_rows, n_cols, sparsity).Seed(3).GenerateDMatrix();
const SparsePage& batch = *dmat->GetBatches<xgboost::SparsePage>().begin();

Expand Down
64 changes: 64 additions & 0 deletions tests/cpp/tree/gpu_hist/test_evaluate_splits.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,70 @@ void TestEvaluateSingleSplit(bool is_categorical) {
parent_sum.GetHess());
}

TEST(GpuHist, Foobar) {
thrust::device_vector<DeviceSplitCandidate> out_splits(1);
GradientPair parent_sum(6.4f, 12.8f);
TrainParam tparam;
std::vector<std::pair<std::string, std::string>> args{
{"max_depth", "1"},
{"max_leaves", "0"},

// Disable all other parameters.
{"colsample_bynode", "1"},
{"colsample_bylevel", "1"},
{"colsample_bytree", "1"},
{"min_child_weight", "0.01"},
{"reg_alpha", "0"},
{"reg_lambda", "0"},
{"max_delta_step", "0"}};
tparam.Init(args);
GPUTrainingParam param{tparam};

thrust::device_vector<bst_feature_t> feature_set =
std::vector<bst_feature_t>{0, 1, 2, 3, 4, 5, 6, 7};
thrust::device_vector<uint32_t> feature_segments =
std::vector<bst_row_t>{0, 3, 6, 9, 12, 15, 18, 21, 24};
thrust::device_vector<float> feature_values =
std::vector<float>{0.30f, 0.67f, 1.64f,
0.32f, 0.77f, 1.95f,
0.29f, 0.70f, 1.80f,
0.32f, 0.75f, 1.85f,
0.18f, 0.59f, 1.69f,
0.25f, 0.74f, 2.00f,
0.26f, 0.74f, 1.98f,
0.26f, 0.71f, 1.83f};
thrust::device_vector<float> feature_min_values =
std::vector<float>{0.1f, 0.2f, 0.3f, 0.1f, 0.2f, 0.3f, 0.2f, 0.2f};
// Setup gradients so that second feature gets higher gain
thrust::device_vector<GradientPair> feature_histogram =
std::vector<GradientPair>{
{0.8314f, 0.7147f}, {1.7989f, 3.7312f}, {3.3846f, 3.4598f},
{2.9277f, 3.5886f}, {1.8429f, 2.4152f}, {1.2443f, 1.9019f},
{1.6380f, 2.9174f}, {1.5657f, 2.5107f}, {2.8111f, 2.4776f},
{2.1322f, 3.0651f}, {3.2927f, 3.8540f}, {0.5899f, 0.9866f},
{1.5185f, 1.6263f}, {2.0686f, 3.1844f}, {2.4278f, 3.0950f},
{1.5105f, 2.1403f}, {2.6922f, 4.2217f}, {1.8122f, 1.5437f},
{0.0000f, 0.0000f}, {4.3245f, 5.7955f}, {1.6903f, 2.1103f},
{2.4012f, 4.4754f}, {3.6136f, 3.4303f}, {0.0000f, 0.0000f}};

common::Span<FeatureType> d_feature_types;
EvaluateSplitInputs<GradientPair> input{1,
parent_sum,
param,
dh::ToSpan(feature_set),
d_feature_types,
dh::ToSpan(feature_segments),
dh::ToSpan(feature_values),
dh::ToSpan(feature_min_values),
dh::ToSpan(feature_histogram)};
TreeEvaluator tree_evaluator(tparam, feature_min_values.size(), 0);
auto evaluator = tree_evaluator.GetEvaluator<GPUTrainingParam>();
EvaluateSingleSplit(dh::ToSpan(out_splits), evaluator, input);

DeviceSplitCandidate result = out_splits[0];
LOG(CONSOLE) << "findex = " << result.findex << ", fvalue = " << result.fvalue;
}

TEST(GpuHist, EvaluateSingleSplit) {
TestEvaluateSingleSplit(false);
}
Expand Down
2 changes: 1 addition & 1 deletion tests/cpp/tree/test_gpu_hist.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -257,7 +257,7 @@ TEST(GpuHist, EvaluateRootSplit) {

DeviceSplitCandidate res = maker.EvaluateRootSplit({6.4f, 12.8f});

ASSERT_EQ(res.findex, 7);
EXPECT_EQ(res.findex, 7);
ASSERT_NEAR(res.fvalue, 0.26, xgboost::kRtEps);
}

Expand Down

0 comments on commit b32a634

Please sign in to comment.