Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

RF node queue rewrite #4125

Merged
merged 21 commits into from
Aug 13, 2021
Merged

RF node queue rewrite #4125

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
812d9c7
Rf gtest rewrite
RAMitchell Jul 8, 2021
2ac77c3
Merge branch 'branch-21.08' of https://github.com/rapidsai/cuml into …
RAMitchell Jul 8, 2021
735c983
Use c++17 constexpr if
RAMitchell Jul 9, 2021
b077cf3
Merge branch 'branch-21.08' of https://github.com/rapidsai/cuml into …
RAMitchell Jul 11, 2021
e54fb6d
Constructor
RAMitchell Jul 12, 2021
1f88053
Lint
RAMitchell Jul 12, 2021
af0aede
Fix uninitialised
RAMitchell Jul 15, 2021
8e57f41
Merge branch 'tests' of github.com:RAMitchell/cuml into shmem-bins
RAMitchell Jul 15, 2021
3c9a27b
Merge branch 'branch-21.08' of https://github.com/rapidsai/cuml into …
RAMitchell Jul 20, 2021
14cb0f5
Refactor
RAMitchell Jul 25, 2021
26362b0
Predict after
RAMitchell Jul 27, 2021
bc0934f
Node queue with work items
RAMitchell Jul 29, 2021
f288814
Fix tests
RAMitchell Aug 1, 2021
c7102ea
Merge branch 'branch-21.10' of https://github.com/rapidsai/cuml into …
RAMitchell Aug 1, 2021
9a0edb9
Relax accuracy test
RAMitchell Aug 2, 2021
6328fe8
Include checker
RAMitchell Aug 2, 2021
ce3221c
Node instance counts
RAMitchell Aug 3, 2021
5019199
Address review comments
RAMitchell Aug 9, 2021
919244b
Batch leaf calculations
RAMitchell Aug 11, 2021
62be8fe
Merge branch 'branch-21.10' of https://github.com/rapidsai/cuml into …
RAMitchell Aug 11, 2021
74b3bbe
Fix warnings
RAMitchell Aug 12, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 6 additions & 7 deletions cpp/include/cuml/tree/flatnode.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,13 +26,13 @@
*/
template <typename DataT, typename LabelT, typename IdxT = int>
struct SparseTreeNode {
LabelT prediction;
IdxT colid = IdxT(-1);
DataT quesval;
DataT best_metric_val;
LabelT prediction = LabelT(0);
IdxT colid = IdxT(-1);
DataT quesval = DataT(0);
DataT best_metric_val = DataT(0);
IdxT left_child_id = IdxT(-1);
uint32_t unique_id = UINT32_MAX;
uint32_t instance_count = UINT32_MAX; // UINT32_MAX indicates n/a
bool IsLeaf() { return left_child_id == -1; }
};

template <typename DataT, typename LabelT, typename IdxT>
Expand All @@ -41,8 +41,7 @@ bool operator==(const SparseTreeNode<DataT, LabelT, IdxT>& lhs,
{
return (lhs.prediction == rhs.prediction) && (lhs.colid == rhs.colid) &&
(lhs.quesval == rhs.quesval) && (lhs.best_metric_val == rhs.best_metric_val) &&
(lhs.left_child_id == rhs.left_child_id) && (lhs.unique_id == rhs.unique_id) &&
(lhs.instance_count == rhs.instance_count);
(lhs.left_child_id == rhs.left_child_id) && (lhs.instance_count == rhs.instance_count);
}

template <typename T, typename L>
Expand Down
81 changes: 27 additions & 54 deletions cpp/src/decisiontree/batched-levelalgo/builder.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,10 +16,7 @@

#pragma once

#include <cuml/common/device_buffer.hpp>
#include <cuml/common/host_buffer.hpp>
#include <raft/mr/device/allocator.hpp>
#include <raft/mr/host/allocator.hpp>
#include <raft/handle.hpp>

#include "builder_base.cuh"
#include "metrics.cuh"
Expand All @@ -33,16 +30,14 @@ template <typename ObjectiveT,
typename DataT = typename ObjectiveT::DataT,
typename LabelT = typename ObjectiveT::LabelT,
typename IdxT = typename ObjectiveT::IdxT>
void convertToSparse(const Builder<ObjectiveT>& b,
const Node<DataT, LabelT, IdxT>* h_nodes,
void convertToSparse(const std::vector<Node<DataT, LabelT, IdxT>> h_nodes,
std::vector<SparseTreeNode<DataT, LabelT>>& sparsetree)
{
auto len = sparsetree.size();
sparsetree.resize(len + b.h_total_nodes);
for (IdxT i = 0; i < b.h_total_nodes; ++i) {
const auto& hnode = h_nodes[i].info;
sparsetree[i + len] = hnode;
sparsetree[i + len].instance_count = h_nodes[i].count;
sparsetree.resize(len + h_nodes.size());
for (std::size_t i = 0; i < h_nodes.size(); ++i) {
const auto& hnode = h_nodes[i].info;
sparsetree[i + len] = hnode;
if (hnode.left_child_id != -1) sparsetree[i + len].left_child_id += len;
}
}
Expand All @@ -52,8 +47,7 @@ template <typename ObjectiveT,
typename DataT = typename ObjectiveT::DataT,
typename LabelT = typename ObjectiveT::LabelT,
typename IdxT = typename ObjectiveT::IdxT>
void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
std::shared_ptr<raft::mr::host::allocator> h_allocator,
void grow_tree(const raft::handle_t& handle,
const DataT* data,
IdxT treeid,
uint64_t seed,
Expand All @@ -65,39 +59,27 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
int n_sampled_rows,
int unique_labels,
const DecisionTreeParams& params,
cudaStream_t stream,
std::vector<SparseTreeNode<DataT, LabelT>>& sparsetree,
IdxT& num_leaves,
IdxT& depth)
{
ML::PUSH_RANGE("DT::grow_tree in batched-levelalgo @builder.cuh");
Builder<ObjectiveT> builder;
size_t d_wsize, h_wsize;
builder.workspaceSize(d_wsize,
h_wsize,
treeid,
seed,
params,
data,
labels,
nrows,
ncols,
n_sampled_rows,
IdxT(params.max_features * ncols),
rowids,
unique_labels,
quantiles);
MLCommon::device_buffer<char> d_buff(d_allocator, stream, d_wsize);
MLCommon::host_buffer<char> h_buff(h_allocator, stream, h_wsize);

std::vector<Node<DataT, LabelT, IdxT>> h_nodes;
h_nodes.reserve(builder.maxNodes);
builder.assignWorkspace(d_buff.data(), h_buff.data());
builder.train(h_nodes, num_leaves, depth, stream);
CUDA_CHECK(cudaStreamSynchronize(stream));
d_buff.release(stream);
h_buff.release(stream);
convertToSparse<ObjectiveT>(builder, h_nodes.data(), sparsetree);
Builder<ObjectiveT> builder(handle,
treeid,
seed,
params,
data,
labels,
nrows,
ncols,
n_sampled_rows,
rowids,
unique_labels,
quantiles);
auto h_nodes = builder.train(num_leaves, depth, handle.get_stream());
CUDA_CHECK(cudaStreamSynchronize(handle.get_stream()));
convertToSparse<ObjectiveT>(h_nodes, sparsetree);
ML::POP_RANGE();
}

Expand All @@ -109,8 +91,7 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
* @tparam LabelT label type
* @tparam IdxT index type
*
* @param[in] d_allocator device allocator
* @param[in] h_allocator host allocator
* @param[in] handle raft handle
* @param[in] data input dataset [on device] [col-major]
* [dim = nrows x ncols]
* @param[in] ncols number of features in the dataset
Expand All @@ -133,8 +114,7 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
* @{
*/
template <typename DataT, typename LabelT, typename IdxT>
void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
std::shared_ptr<raft::mr::host::allocator> h_allocator,
void grow_tree(const raft::handle_t& handle,
const DataT* data,
IdxT treeid,
uint64_t seed,
Expand All @@ -146,15 +126,13 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
int n_sampled_rows,
int unique_labels,
const DecisionTreeParams& params,
cudaStream_t stream,
std::vector<SparseTreeNode<DataT, LabelT>>& sparsetree,
IdxT& num_leaves,
IdxT& depth)
{
// Dispatch objective
if (params.split_criterion == CRITERION::GINI) {
grow_tree<GiniObjectiveFunction<DataT, LabelT, IdxT>>(d_allocator,
h_allocator,
grow_tree<GiniObjectiveFunction<DataT, LabelT, IdxT>>(handle,
data,
treeid,
seed,
Expand All @@ -166,13 +144,11 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
n_sampled_rows,
unique_labels,
params,
stream,
sparsetree,
num_leaves,
depth);
} else if (params.split_criterion == CRITERION::ENTROPY) {
grow_tree<EntropyObjectiveFunction<DataT, LabelT, IdxT>>(d_allocator,
h_allocator,
grow_tree<EntropyObjectiveFunction<DataT, LabelT, IdxT>>(handle,
data,
treeid,
seed,
Expand All @@ -184,13 +160,11 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
n_sampled_rows,
unique_labels,
params,
stream,
sparsetree,
num_leaves,
depth);
} else if (params.split_criterion == CRITERION::MSE) {
grow_tree<MSEObjectiveFunction<DataT, LabelT, IdxT>>(d_allocator,
h_allocator,
grow_tree<MSEObjectiveFunction<DataT, LabelT, IdxT>>(handle,
data,
treeid,
seed,
Expand All @@ -202,7 +176,6 @@ void grow_tree(std::shared_ptr<raft::mr::device::allocator> d_allocator,
n_sampled_rows,
unique_labels,
params,
stream,
sparsetree,
num_leaves,
depth);
Expand Down
Loading