GPU multiclass metrics

src/metric/elementwise_metric.cu

@@ -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
@@ -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
@@ -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(
@@ -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)
@@ -370,7 +349,7 @@ struct EvalEWiseBase : public Metric {
 
   MetricParam param_;
 
-  MetricsReduction<Policy> reducer_;
+  ElementWiseMetricsReduction<Policy> reducer_;
 };
 
 XGBOOST_REGISTER_METRIC(RMSE, "rmse")

src/metric/metric.cc

@@ -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);

src/metric/metric_common.h

@@ -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_

src/metric/multiclass_metric.cc

@@ -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)