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

Fix GPU ID and prediction cache from pickle (#5086)

Browse Source 
* Hack for saving GPU ID.

* Declare prediction cache on GBTree.

* Add a simple test.

* Add `auto` option for GPU Predictor.
This commit is contained in:
Jiaming Yuan
2019-12-07 16:02:06 +08:00
committed by GitHub
parent 7ef5b78003
commit 608ebbe444
17 changed files with 362 additions and 182 deletions
Download Patch File Download Diff File Expand all files Collapse all files

72
src/predictor/cpu_predictor.cc
View File

@@ -1,6 +1,8 @@
/*!
* Copyright by Contributors 2017
* Copyright by Contributors 2017-2019
*/
#include <dmlc/omp.h>
#include "xgboost/predictor.h"
#include "xgboost/tree_model.h"
#include "xgboost/tree_updater.h"
@@ -43,10 +45,11 @@ class CPUPredictor : public Predictor {
}
}
}
inline void PredLoopSpecalize(DMatrix* p_fmat,
std::vector<bst_float>* out_preds,
const gbm::GBTreeModel& model, int num_group,
unsigned tree_begin, unsigned tree_end) {
void PredLoopInternal(DMatrix* p_fmat, std::vector<bst_float>* out_preds,
gbm::GBTreeModel const& model, int32_t tree_begin,
int32_t tree_end) {
int32_t const num_group = model.param.num_output_group;
const int nthread = omp_get_max_threads();
InitThreadTemp(nthread, model.param.num_feature);
std::vector<bst_float>& preds = *out_preds;
@@ -99,22 +102,15 @@ class CPUPredictor : public Predictor {
}
}
void PredLoopInternal(DMatrix* dmat, std::vector<bst_float>* out_preds,
const gbm::GBTreeModel& model, int tree_begin,
unsigned ntree_limit) {
// TODO(Rory): Check if this specialisation actually improves performance
PredLoopSpecalize(dmat, out_preds, model, model.param.num_output_group,
tree_begin, ntree_limit);
}
bool PredictFromCache(DMatrix* dmat,
HostDeviceVector<bst_float>* out_preds,
const gbm::GBTreeModel& model,
unsigned ntree_limit) {
unsigned ntree_limit) const {
CHECK(cache_);
if (ntree_limit == 0 ||
ntree_limit * model.param.num_output_group >= model.trees.size()) {
auto it = cache_.find(dmat);
if (it != cache_.end()) {
auto it = cache_->find(dmat);
if (it != cache_->end()) {
const HostDeviceVector<bst_float>& y = it->second.predictions;
if (y.Size() != 0) {
out_preds->Resize(y.Size());
@@ -130,6 +126,7 @@ class CPUPredictor : public Predictor {
void InitOutPredictions(const MetaInfo& info,
HostDeviceVector<bst_float>* out_preds,
const gbm::GBTreeModel& model) const {
CHECK_NE(model.param.num_output_group, 0);
size_t n = model.param.num_output_group * info.num_row_;
const auto& base_margin = info.base_margin_.HostVector();
out_preds->Resize(n);
@@ -150,21 +147,24 @@ class CPUPredictor : public Predictor {
oss << "[number of data points], i.e. " << info.num_row_ << ". ";
}
oss << "Instead, all data points will use "
<< "base_score = " << model.base_margin;
<< "base_margin = " << model.base_margin;
LOG(WARNING) << oss.str();
}
std::fill(out_preds_h.begin(), out_preds_h.end(), model.base_margin);
std::fill(out_preds_h.begin(), out_preds_h.end(),
model.base_margin);
}
}
public:
CPUPredictor(GenericParameter const* generic_param,
std::shared_ptr<std::unordered_map<DMatrix*, PredictionCacheEntry>> cache) :
Predictor::Predictor{generic_param, cache} {}
void PredictBatch(DMatrix* dmat, HostDeviceVector<bst_float>* out_preds,
const gbm::GBTreeModel& model, int tree_begin,
unsigned ntree_limit = 0) override {
if (this->PredictFromCache(dmat, out_preds, model, ntree_limit)) {
return;
}
this->InitOutPredictions(dmat->Info(), out_preds, model);
ntree_limit *= model.param.num_output_group;
@@ -174,6 +174,15 @@ class CPUPredictor : public Predictor {
this->PredLoopInternal(dmat, &out_preds->HostVector(), model,
tree_begin, ntree_limit);
auto cache_emtry = this->FindCache(dmat);
if (cache_emtry == cache_->cend()) { return; }
if (cache_emtry->second.predictions.Size() == 0) {
// See comment in GPUPredictor::PredictBatch.
InitOutPredictions(cache_emtry->second.data->Info(),
&(cache_emtry->second.predictions), model);
cache_emtry->second.predictions.Copy(*out_preds);
}
}
void UpdatePredictionCache(
@@ -182,7 +191,7 @@ class CPUPredictor : public Predictor {
int num_new_trees) override {
int old_ntree = model.trees.size() - num_new_trees;
// update cache entry
for (auto& kv : cache_) {
for (auto& kv : (*cache_)) {
PredictionCacheEntry& e = kv.second;
if (e.predictions.Size() == 0) {
@@ -215,7 +224,7 @@ class CPUPredictor : public Predictor {
out_preds->resize(model.param.num_output_group *
(model.param.size_leaf_vector + 1));
// loop over output groups
for (int gid = 0; gid < model.param.num_output_group; ++gid) {
for (uint32_t gid = 0; gid < model.param.num_output_group; ++gid) {
(*out_preds)[gid] =
PredValue(inst, model.trees, model.tree_info, gid,
&thread_temp[0], 0, ntree_limit) +
@@ -254,10 +263,9 @@ class CPUPredictor : public Predictor {
}
void PredictContribution(DMatrix* p_fmat, std::vector<bst_float>* out_contribs,
const gbm::GBTreeModel& model, unsigned ntree_limit,
const gbm::GBTreeModel& model, uint32_t ntree_limit,
std::vector<bst_float>* tree_weights,
bool approximate,
int condition,
bool approximate, int condition,
unsigned condition_feature) override {
const int nthread = omp_get_max_threads();
InitThreadTemp(nthread, model.param.num_feature);
@@ -268,7 +276,9 @@ class CPUPredictor : public Predictor {
ntree_limit = static_cast<unsigned>(model.trees.size());
}
const int ngroup = model.param.num_output_group;
CHECK_NE(ngroup, 0);
size_t const ncolumns = model.param.num_feature + 1;
CHECK_NE(ncolumns, 0);
// allocate space for (number of features + bias) times the number of rows
std::vector<bst_float>& contribs = *out_contribs;
contribs.resize(info.num_row_ * ncolumns * model.param.num_output_group);
@@ -292,8 +302,7 @@ class CPUPredictor : public Predictor {
RegTree::FVec& feats = thread_temp[omp_get_thread_num()];
// loop over all classes
for (int gid = 0; gid < ngroup; ++gid) {
bst_float* p_contribs =
&contribs[(row_idx * ngroup + gid) * ncolumns];
bst_float* p_contribs = &contribs[(row_idx * ngroup + gid) * ncolumns];
feats.Fill(batch[i]);
// calculate contributions
for (unsigned j = 0; j < ntree_limit; ++j) {
@@ -307,7 +316,7 @@ class CPUPredictor : public Predictor {
} else {
model.trees[j]->CalculateContributionsApprox(feats, &this_tree_contribs[0]);
}
for (int ci = 0 ; ci < ncolumns ; ++ci) {
for (size_t ci = 0 ; ci < ncolumns ; ++ci) {
p_contribs[ci] += this_tree_contribs[ci] *
(tree_weights == nullptr ? 1 : (*tree_weights)[j]);
}
@@ -330,7 +339,7 @@ class CPUPredictor : public Predictor {
bool approximate) override {
const MetaInfo& info = p_fmat->Info();
const int ngroup = model.param.num_output_group;
size_t ncolumns = model.param.num_feature;
size_t const ncolumns = model.param.num_feature;
const unsigned row_chunk = ngroup * (ncolumns + 1) * (ncolumns + 1);
const unsigned mrow_chunk = (ncolumns + 1) * (ncolumns + 1);
const unsigned crow_chunk = ngroup * (ncolumns + 1);
@@ -375,7 +384,10 @@ class CPUPredictor : public Predictor {
};
XGBOOST_REGISTER_PREDICTOR(CPUPredictor, "cpu_predictor")
.describe("Make predictions using CPU.")
.set_body([]() { return new CPUPredictor(); });
.describe("Make predictions using CPU.")
.set_body([](GenericParameter const* generic_param,
std::shared_ptr<std::unordered_map<DMatrix*, PredictionCacheEntry>> cache) {
return new CPUPredictor(generic_param, cache);
});
} // namespace predictor
} // namespace xgboost

87
src/predictor/gpu_predictor.cu
View File

@@ -202,7 +202,7 @@ class GPUPredictor : public xgboost::Predictor {
const thrust::host_vector<size_t>& h_tree_segments,
const thrust::host_vector<DevicePredictionNode>& h_nodes,
size_t tree_begin, size_t tree_end) {
dh::safe_cuda(cudaSetDevice(device_));
dh::safe_cuda(cudaSetDevice(generic_param_->gpu_id));
nodes_.resize(h_nodes.size());
dh::safe_cuda(cudaMemcpyAsync(nodes_.data().get(), h_nodes.data(),
sizeof(DevicePredictionNode) * h_nodes.size(),
@@ -224,7 +224,11 @@ class GPUPredictor : public xgboost::Predictor {
size_t num_features,
HostDeviceVector<bst_float>* predictions,
size_t batch_offset) {
dh::safe_cuda(cudaSetDevice(device_));
dh::safe_cuda(cudaSetDevice(generic_param_->gpu_id));
batch.data.SetDevice(generic_param_->gpu_id);
batch.offset.SetDevice(generic_param_->gpu_id);
predictions->SetDevice(generic_param_->gpu_id);
const uint32_t BLOCK_THREADS = 128;
size_t num_rows = batch.Size();
auto GRID_SIZE = static_cast<uint32_t>(common::DivRoundUp(num_rows, BLOCK_THREADS));
@@ -271,16 +275,19 @@ class GPUPredictor : public xgboost::Predictor {
HostDeviceVector<bst_float>* out_preds,
const gbm::GBTreeModel& model, size_t tree_begin,
size_t tree_end) {
if (tree_end - tree_begin == 0) { return; }
if (tree_end - tree_begin == 0) {
return;
}
monitor_.StartCuda("DevicePredictInternal");
InitModel(model, tree_begin, tree_end);
size_t batch_offset = 0;
for (auto &batch : dmat->GetBatches<SparsePage>()) {
batch.offset.SetDevice(device_);
batch.data.SetDevice(device_);
PredictInternal(batch, model.param.num_feature, out_preds, batch_offset);
batch.offset.SetDevice(generic_param_->gpu_id);
batch.data.SetDevice(generic_param_->gpu_id);
PredictInternal(batch, model.param.num_feature,
out_preds, batch_offset);
batch_offset += batch.Size() * model.param.num_output_group;
}
@@ -288,19 +295,21 @@ class GPUPredictor : public xgboost::Predictor {
}
public:
GPUPredictor() : device_{-1} {}
GPUPredictor(GenericParameter const* generic_param,
std::shared_ptr<std::unordered_map<DMatrix*, PredictionCacheEntry>> cache) :
Predictor::Predictor{generic_param, cache} {}
~GPUPredictor() override {
if (device_ >= 0) {
dh::safe_cuda(cudaSetDevice(device_));
if (generic_param_->gpu_id >= 0) {
dh::safe_cuda(cudaSetDevice(generic_param_->gpu_id));
}
}
void PredictBatch(DMatrix* dmat, HostDeviceVector<bst_float>* out_preds,
const gbm::GBTreeModel& model, int tree_begin,
unsigned ntree_limit = 0) override {
int device = learner_param_->gpu_id;
CHECK_GE(device, 0);
int device = generic_param_->gpu_id;
CHECK_GE(device, 0) << "Set `gpu_id' to positive value for processing GPU data.";
ConfigureDevice(device);
if (this->PredictFromCache(dmat, out_preds, model, ntree_limit)) {
@@ -308,13 +317,30 @@ class GPUPredictor : public xgboost::Predictor {
}
this->InitOutPredictions(dmat->Info(), out_preds, model);
int tree_end = ntree_limit * model.param.num_output_group;
int32_t tree_end = ntree_limit * model.param.num_output_group;
if (ntree_limit == 0 || ntree_limit > model.trees.size()) {
tree_end = static_cast<unsigned>(model.trees.size());
}
DevicePredictInternal(dmat, out_preds, model, tree_begin, tree_end);
auto cache_emtry = this->FindCache(dmat);
if (cache_emtry == cache_->cend()) { return; }
if (cache_emtry->second.predictions.Size() == 0) {
// Initialise the cache on first iteration, this comes useful
// when performing training continuation:
//
// 1. PredictBatch
// 2. CommitModel
// - updater->UpdatePredictionCache
//
// If we don't initialise this cache, the 2 step will recieve an invalid cache as
// the first step only modifies prediction store in learner without following code.
InitOutPredictions(cache_emtry->second.data->Info(),
&(cache_emtry->second.predictions), model);
cache_emtry->second.predictions.Copy(*out_preds);
}
}
protected:
@@ -324,7 +350,7 @@ class GPUPredictor : public xgboost::Predictor {
size_t n_classes = model.param.num_output_group;
size_t n = n_classes * info.num_row_;
const HostDeviceVector<bst_float>& base_margin = info.base_margin_;
out_preds->SetDevice(device_);
out_preds->SetDevice(generic_param_->gpu_id);
out_preds->Resize(n);
if (base_margin.Size() != 0) {
CHECK_EQ(base_margin.Size(), n);
@@ -338,8 +364,8 @@ class GPUPredictor : public xgboost::Predictor {
const gbm::GBTreeModel& model, unsigned ntree_limit) {
if (ntree_limit == 0 ||
ntree_limit * model.param.num_output_group >= model.trees.size()) {
auto it = cache_.find(dmat);
if (it != cache_.end()) {
auto it = (*cache_).find(dmat);
if (it != cache_->cend()) {
const HostDeviceVector<bst_float>& y = it->second.predictions;
if (y.Size() != 0) {
monitor_.StartCuda("PredictFromCache");
@@ -360,7 +386,7 @@ class GPUPredictor : public xgboost::Predictor {
int num_new_trees) override {
auto old_ntree = model.trees.size() - num_new_trees;
// update cache entry
for (auto& kv : cache_) {
for (auto& kv : (*cache_)) {
PredictionCacheEntry& e = kv.second;
DMatrix* dmat = kv.first;
HostDeviceVector<bst_float>& predictions = e.predictions;
@@ -382,14 +408,14 @@ class GPUPredictor : public xgboost::Predictor {
void PredictInstance(const SparsePage::Inst& inst,
std::vector<bst_float>* out_preds,
const gbm::GBTreeModel& model, unsigned ntree_limit) override {
LOG(FATAL) << "Internal error: " << __func__
LOG(FATAL) << "[Internal error]: " << __func__
<< " is not implemented in GPU Predictor.";
}
void PredictLeaf(DMatrix* p_fmat, std::vector<bst_float>* out_preds,
const gbm::GBTreeModel& model,
unsigned ntree_limit) override {
LOG(FATAL) << "Internal error: " << __func__
LOG(FATAL) << "[Internal error]: " << __func__
<< " is not implemented in GPU Predictor.";
}
@@ -399,7 +425,7 @@ class GPUPredictor : public xgboost::Predictor {
std::vector<bst_float>* tree_weights,
bool approximate, int condition,
unsigned condition_feature) override {
LOG(FATAL) << "Internal error: " << __func__
LOG(FATAL) << "[Internal error]: " << __func__
<< " is not implemented in GPU Predictor.";
}
@@ -409,15 +435,14 @@ class GPUPredictor : public xgboost::Predictor {
unsigned ntree_limit,
std::vector<bst_float>* tree_weights,
bool approximate) override {
LOG(FATAL) << "Internal error: " << __func__
LOG(FATAL) << "[Internal error]: " << __func__
<< " is not implemented in GPU Predictor.";
}
void Configure(const std::vector<std::pair<std::string, std::string>>& cfg,
const std::vector<std::shared_ptr<DMatrix>>& cache) override {
Predictor::Configure(cfg, cache);
void Configure(const std::vector<std::pair<std::string, std::string>>& cfg) override {
Predictor::Configure(cfg);
int device = learner_param_->gpu_id;
int device = generic_param_->gpu_id;
if (device >= 0) {
ConfigureDevice(device);
}
@@ -426,14 +451,11 @@ class GPUPredictor : public xgboost::Predictor {
private:
/*! \brief Reconfigure the device when GPU is changed. */
void ConfigureDevice(int device) {
if (device_ == device) return;
device_ = device;
if (device_ >= 0) {
max_shared_memory_bytes_ = dh::MaxSharedMemory(device_);
if (device >= 0) {
max_shared_memory_bytes_ = dh::MaxSharedMemory(device);
}
}
int device_;
common::Monitor monitor_;
dh::device_vector<DevicePredictionNode> nodes_;
dh::device_vector<size_t> tree_segments_;
@@ -445,8 +467,11 @@ class GPUPredictor : public xgboost::Predictor {
};
XGBOOST_REGISTER_PREDICTOR(GPUPredictor, "gpu_predictor")
.describe("Make predictions using GPU.")
.set_body([]() { return new GPUPredictor(); });
.describe("Make predictions using GPU.")
.set_body([](GenericParameter const* generic_param,
std::shared_ptr<std::unordered_map<DMatrix*, PredictionCacheEntry>> cache) {
return new GPUPredictor(generic_param, cache);
});
} // namespace predictor
} // namespace xgboost

13
src/predictor/predictor.cc
View File

@@ -9,19 +9,16 @@ DMLC_REGISTRY_ENABLE(::xgboost::PredictorReg);
} // namespace dmlc
namespace xgboost {
void Predictor::Configure(
const std::vector<std::pair<std::string, std::string>>& cfg,
const std::vector<std::shared_ptr<DMatrix>>& cache) {
for (const std::shared_ptr<DMatrix>& d : cache) {
cache_[d.get()].data = d;
}
const std::vector<std::pair<std::string, std::string>>& cfg) {
}
Predictor* Predictor::Create(std::string const& name, GenericParameter const* learner_param) {
Predictor* Predictor::Create(
std::string const& name, GenericParameter const* generic_param,
std::shared_ptr<std::unordered_map<DMatrix*, PredictionCacheEntry>> cache) {
auto* e = ::dmlc::Registry<PredictorReg>::Get()->Find(name);
if (e == nullptr) {
LOG(FATAL) << "Unknown predictor type " << name;
}
auto p_predictor = (e->body)();
p_predictor->learner_param_ = learner_param;
auto p_predictor = (e->body)(generic_param, cache);
return p_predictor;
}
} // namespace xgboost