Skip to content

Commit

Permalink
GPU multiclass metrics (#4368)
Browse files Browse the repository at this point in the history
* Port multi classes metrics to CUDA.
  • Loading branch information
trivialfis authored Apr 15, 2019
1 parent be7bc07 commit 84d992b
Show file tree
Hide file tree
Showing 8 changed files with 368 additions and 189 deletions.
41 changes: 10 additions & 31 deletions src/metric/elementwise_metric.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*!
* Copyright 2015-2018 by Contributors
* Copyright 2015-2019 by Contributors
* \file elementwise_metric.cc
* \brief evaluation metrics for elementwise binary or regression.
* \author Kailong Chen, Tianqi Chen
Expand All @@ -9,15 +9,15 @@
#include <dmlc/registry.h>
#include <cmath>

#include "metric_param.h"
#include "metric_common.h"
#include "../common/math.h"
#include "../common/common.h"

#if defined(XGBOOST_USE_CUDA)
#include <thrust/iterator/counting_iterator.h>
#include <thrust/execution_policy.h> // thrust::cuda::par
#include <thrust/functional.h> // thrust::plus<>
#include <thrust/transform_reduce.h>
#include <thrust/execution_policy.h>
#include <thrust/functional.h> // thrust::plus<>
#include <thrust/iterator/counting_iterator.h>

#include "../common/device_helpers.cuh"
#endif // XGBOOST_USE_CUDA
Expand All @@ -28,29 +28,9 @@ namespace metric {
DMLC_REGISTRY_FILE_TAG(elementwise_metric);

template <typename EvalRow>
class MetricsReduction {
public:
class PackedReduceResult {
double residue_sum_;
double weights_sum_;
friend MetricsReduction;

public:
XGBOOST_DEVICE PackedReduceResult() : residue_sum_{0}, weights_sum_{0} {}
XGBOOST_DEVICE PackedReduceResult(double residue, double weight) :
residue_sum_{residue}, weights_sum_{weight} {}

XGBOOST_DEVICE
PackedReduceResult operator+(PackedReduceResult const& other) const {
return PackedReduceResult { residue_sum_ + other.residue_sum_,
weights_sum_ + other.weights_sum_ };
}
double Residue() const { return residue_sum_; }
double Weights() const { return weights_sum_; }
};

class ElementWiseMetricsReduction {
public:
explicit MetricsReduction(EvalRow policy) :
explicit ElementWiseMetricsReduction(EvalRow policy) :
policy_(std::move(policy)) {}

PackedReduceResult CpuReduceMetrics(
Expand Down Expand Up @@ -144,9 +124,8 @@ class MetricsReduction {
DeviceReduceMetrics(id, index, weights, labels, preds);
}

for (size_t i = 0; i < devices.Size(); ++i) {
result.residue_sum_ += res_per_device[i].residue_sum_;
result.weights_sum_ += res_per_device[i].weights_sum_;
for (auto const& res : res_per_device) {
result += res;
}
}
#endif // defined(XGBOOST_USE_CUDA)
Expand Down Expand Up @@ -370,7 +349,7 @@ struct EvalEWiseBase : public Metric {

MetricParam param_;

MetricsReduction<Policy> reducer_;
ElementWiseMetricsReduction<Policy> reducer_;
};

XGBOOST_REGISTER_METRIC(RMSE, "rmse")
Expand Down
2 changes: 1 addition & 1 deletion src/metric/metric.cc
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
#include <xgboost/metric.h>
#include <dmlc/registry.h>

#include "metric_param.h"
#include "metric_common.h"

namespace dmlc {
DMLC_REGISTRY_ENABLE(::xgboost::MetricReg);
Expand Down
54 changes: 54 additions & 0 deletions src/metric/metric_common.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,54 @@
/*!
* Copyright 2018-2019 by Contributors
* \file metric_param.cc
*/
#ifndef XGBOOST_METRIC_METRIC_COMMON_H_
#define XGBOOST_METRIC_METRIC_COMMON_H_

#include <dmlc/parameter.h>
#include "../common/common.h"

namespace xgboost {
namespace metric {

// Created exclusively for GPU.
struct MetricParam : public dmlc::Parameter<MetricParam> {
int n_gpus;
int gpu_id;
DMLC_DECLARE_PARAMETER(MetricParam) {
DMLC_DECLARE_FIELD(n_gpus).set_default(1).set_lower_bound(GPUSet::kAll)
.describe("Number of GPUs to use for multi-gpu algorithms.");
DMLC_DECLARE_FIELD(gpu_id)
.set_lower_bound(0)
.set_default(0)
.describe("gpu to use for objective function evaluation");
};
};

class PackedReduceResult {
double residue_sum_;
double weights_sum_;

public:
XGBOOST_DEVICE PackedReduceResult() : residue_sum_{0}, weights_sum_{0} {}
XGBOOST_DEVICE PackedReduceResult(double residue, double weight)
: residue_sum_{residue}, weights_sum_{weight} {}

XGBOOST_DEVICE
PackedReduceResult operator+(PackedReduceResult const &other) const {
return PackedReduceResult{residue_sum_ + other.residue_sum_,
weights_sum_ + other.weights_sum_};
}
PackedReduceResult &operator+=(PackedReduceResult const &other) {
this->residue_sum_ += other.residue_sum_;
this->weights_sum_ += other.weights_sum_;
return *this;
}
double Residue() const { return residue_sum_; }
double Weights() const { return weights_sum_; }
};

} // namespace metric
} // namespace xgboost

#endif // XGBOOST_METRIC_METRIC_COMMON_H_
31 changes: 0 additions & 31 deletions src/metric/metric_param.h

This file was deleted.

128 changes: 5 additions & 123 deletions src/metric/multiclass_metric.cc
Original file line number Diff line number Diff line change
@@ -1,126 +1,8 @@
/*!
* Copyright 2015 by Contributors
* \file multiclass_metric.cc
* \brief evaluation metrics for multiclass classification.
* \author Kailong Chen, Tianqi Chen
* Copyright 2019 XGBoost contributors
*/
#include <rabit/rabit.h>
#include <xgboost/metric.h>
#include <cmath>
#include "../common/math.h"
// Dummy file to keep the CUDA conditional compile trick.

namespace xgboost {
namespace metric {
// tag the this file, used by force static link later.
DMLC_REGISTRY_FILE_TAG(multiclass_metric);

/*!
* \brief base class of multi-class evaluation
* \tparam Derived the name of subclass
*/
template<typename Derived>
struct EvalMClassBase : public Metric {
bst_float Eval(const HostDeviceVector<bst_float> &preds,
const MetaInfo &info,
bool distributed) override {
CHECK_NE(info.labels_.Size(), 0U) << "label set cannot be empty";
CHECK(preds.Size() % info.labels_.Size() == 0)
<< "label and prediction size not match";
const size_t nclass = preds.Size() / info.labels_.Size();
CHECK_GE(nclass, 1U)
<< "mlogloss and merror are only used for multi-class classification,"
<< " use logloss for binary classification";
const auto ndata = static_cast<bst_omp_uint>(info.labels_.Size());
double sum = 0.0, wsum = 0.0;
int label_error = 0;

const auto& labels = info.labels_.HostVector();
const auto& weights = info.weights_.HostVector();
const std::vector<bst_float>& h_preds = preds.HostVector();

#pragma omp parallel for reduction(+: sum, wsum) schedule(static)
for (bst_omp_uint i = 0; i < ndata; ++i) {
const bst_float wt = weights.size() > 0 ? weights[i] : 1.0f;
auto label = static_cast<int>(labels[i]);
if (label >= 0 && label < static_cast<int>(nclass)) {
sum += Derived::EvalRow(label,
h_preds.data() + i * nclass,
nclass) * wt;
wsum += wt;
} else {
label_error = label;
}
}
CHECK(label_error >= 0 && label_error < static_cast<int>(nclass))
<< "MultiClassEvaluation: label must be in [0, num_class),"
<< " num_class=" << nclass << " but found " << label_error << " in label";

double dat[2]; dat[0] = sum, dat[1] = wsum;
if (distributed) {
rabit::Allreduce<rabit::op::Sum>(dat, 2);
}
return Derived::GetFinal(dat[0], dat[1]);
}
/*!
* \brief to be implemented by subclass,
* get evaluation result from one row
* \param label label of current instance
* \param pred prediction value of current instance
* \param nclass number of class in the prediction
*/
inline static bst_float EvalRow(int label,
const bst_float *pred,
size_t nclass);
/*!
* \brief to be overridden by subclass, final transformation
* \param esum the sum statistics returned by EvalRow
* \param wsum sum of weight
*/
inline static bst_float GetFinal(bst_float esum, bst_float wsum) {
return esum / wsum;
}

private:
// used to store error message
const char *error_msg_;
};

/*! \brief match error */
struct EvalMatchError : public EvalMClassBase<EvalMatchError> {
const char* Name() const override {
return "merror";
}
inline static bst_float EvalRow(int label,
const bst_float *pred,
size_t nclass) {
return common::FindMaxIndex(pred, pred + nclass) != pred + static_cast<int>(label);
}
};

/*! \brief match error */
struct EvalMultiLogLoss : public EvalMClassBase<EvalMultiLogLoss> {
const char* Name() const override {
return "mlogloss";
}
inline static bst_float EvalRow(int label,
const bst_float *pred,
size_t nclass) {
const bst_float eps = 1e-16f;
auto k = static_cast<size_t>(label);
if (pred[k] > eps) {
return -std::log(pred[k]);
} else {
return -std::log(eps);
}
}
};

XGBOOST_REGISTER_METRIC(MatchError, "merror")
.describe("Multiclass classification error.")
.set_body([](const char* param) { return new EvalMatchError(); });

XGBOOST_REGISTER_METRIC(MultiLogLoss, "mlogloss")
.describe("Multiclass negative loglikelihood.")
.set_body([](const char* param) { return new EvalMultiLogLoss(); });
} // namespace metric
} // namespace xgboost
#if !defined(XGBOOST_USE_CUDA)
#include "multiclass_metric.cu"
#endif // !defined(XGBOOST_USE_CUDA)
Loading

0 comments on commit 84d992b

Please sign in to comment.