[CUDA] Add binary logloss metric for new CUDA version (#5635)

src/metric/binary_metric.hpp

@@ -96,7 +96,7 @@ class BinaryMetric: public Metric {
     return std::vector<double>(1, loss);
   }
 
- private:
+ protected:
   /*! \brief Number of data */
   data_size_t num_data_;
   /*! \brief Pointer of label */

src/metric/cuda/cuda_binary_metric.cpp

+/*!
+ * Copyright (c) 2022 Microsoft Corporation. All rights reserved.
+ * Licensed under the MIT License. See LICENSE file in the project root for
+ * license information.
+ */
+
+#ifdef USE_CUDA_EXP
+
+#include "cuda_binary_metric.hpp"
+
+namespace LightGBM {
+
+CUDABinaryLoglossMetric::CUDABinaryLoglossMetric(const Config& config):
+  CUDABinaryMetricInterface<BinaryLoglossMetric, CUDABinaryLoglossMetric>(config) {}
+
+template <typename HOST_METRIC, typename CUDA_METRIC>
+std::vector<double> CUDABinaryMetricInterface<HOST_METRIC, CUDA_METRIC>::Eval(const double* score, const ObjectiveFunction* objective) const {
+  const double* score_convert = score;
+  if (objective != nullptr && objective->NeedConvertOutputCUDA()) {
+    this->score_convert_buffer_.Resize(static_cast<size_t>(this->num_data_) * static_cast<size_t>(this->num_class_));
+    score_convert = objective->ConvertOutputCUDA(this->num_data_, score, this->score_convert_buffer_.RawData());
+  }
+  double sum_loss = 0.0, sum_weight = 0.0;
+  this->LaunchEvalKernel(score_convert, &sum_loss, &sum_weight);
+  const double eval_score = sum_loss / sum_weight;
+  return std::vector<double>{eval_score};
+}
+
+}  // namespace LightGBM
+
+#endif  // USE_CUDA_EXP

src/metric/cuda/cuda_binary_metric.hpp

+/*!
+ * Copyright (c) 2022 Microsoft Corporation. All rights reserved.
+ * Licensed under the MIT License. See LICENSE file in the project root for
+ * license information.
+ */
+
+#ifndef LIGHTGBM_METRIC_CUDA_CUDA_BINARY_METRIC_HPP_
+#define LIGHTGBM_METRIC_CUDA_CUDA_BINARY_METRIC_HPP_
+
+#ifdef USE_CUDA_EXP
+
+#include <LightGBM/cuda/cuda_metric.hpp>
+#include <LightGBM/cuda/cuda_utils.h>
+
+#include <vector>
+
+#include "cuda_regression_metric.hpp"
+#include "../binary_metric.hpp"
+
+namespace LightGBM {
+
+template <typename HOST_METRIC, typename CUDA_METRIC>
+class CUDABinaryMetricInterface: public CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC> {
+ public:
+  explicit CUDABinaryMetricInterface(const Config& config): CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>(config) {}
+
+  virtual ~CUDABinaryMetricInterface() {}
+
+  std::vector<double> Eval(const double* score, const ObjectiveFunction* objective) const override;
+};
+
+class CUDABinaryLoglossMetric: public CUDABinaryMetricInterface<BinaryLoglossMetric, CUDABinaryLoglossMetric> {
+ public:
+  explicit CUDABinaryLoglossMetric(const Config& config);
+
+  virtual ~CUDABinaryLoglossMetric() {}
+
+  __device__ static double MetricOnPointCUDA(label_t label, double score) {
+    // score should have been converted to probability
+    if (label <= 0) {
+      if (1.0f - score > kEpsilon) {
+        return -log(1.0f - score);
+      }
+    } else {
+      if (score > kEpsilon) {
+        return -log(score);
+      }
+    }
+    return -log(kEpsilon);
+  }
+};
+
+}  // namespace LightGBM
+
+#endif  // USE_CUDA_EXP
+
+#endif  // LIGHTGBM_METRIC_CUDA_CUDA_BINARY_METRIC_HPP_

src/metric/cuda/cuda_pointwise_metric.cpp

+/*!
+ * Copyright (c) 2022 Microsoft Corporation. All rights reserved.
+ * Licensed under the MIT License. See LICENSE file in the project root for
+ * license information.
+ */
+
+#ifdef USE_CUDA_EXP
+
+#include "cuda_binary_metric.hpp"
+#include "cuda_pointwise_metric.hpp"
+#include "cuda_regression_metric.hpp"
+
+namespace LightGBM {
+
+template <typename HOST_METRIC, typename CUDA_METRIC>
+void CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>::Init(const Metadata& metadata, data_size_t num_data) {
+  CUDAMetricInterface<HOST_METRIC>::Init(metadata, num_data);
+  const int max_num_reduce_blocks = (this->num_data_ + NUM_DATA_PER_EVAL_THREAD - 1) / NUM_DATA_PER_EVAL_THREAD;
+  if (this->cuda_weights_ == nullptr) {
+    reduce_block_buffer_.Resize(max_num_reduce_blocks);
+  } else {
+    reduce_block_buffer_.Resize(max_num_reduce_blocks * 2);
+  }
+  const int max_num_reduce_blocks_inner = (max_num_reduce_blocks + NUM_DATA_PER_EVAL_THREAD - 1) / NUM_DATA_PER_EVAL_THREAD;
+  if (this->cuda_weights_ == nullptr) {
+    reduce_block_buffer_inner_.Resize(max_num_reduce_blocks_inner);
+  } else {
+    reduce_block_buffer_inner_.Resize(max_num_reduce_blocks_inner * 2);
+  }
+}
+
+template void CUDAPointwiseMetricInterface<RMSEMetric, CUDARMSEMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<L2Metric, CUDAL2Metric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<BinaryLoglossMetric, CUDABinaryLoglossMetric>::Init(const Metadata& metadata, data_size_t num_data);
+
+}  // namespace LightGBM
+
+#endif  // USE_CUDA_EXP

src/metric/cuda/cuda_regression_metric.cu → src/metric/cuda/cuda_pointwise_metric.cu

@@ -8,6 +8,8 @@
 
 #include <LightGBM/cuda/cuda_algorithms.hpp>
 
+#include "cuda_binary_metric.hpp"
+#include "cuda_pointwise_metric.hpp"
 #include "cuda_regression_metric.hpp"
 
 namespace LightGBM {
@@ -40,7 +42,7 @@ __global__ void EvalKernel(const data_size_t num_data, const label_t* labels, co
 }
 
 template <typename HOST_METRIC, typename CUDA_METRIC>
-double CUDARegressionMetricInterface<HOST_METRIC, CUDA_METRIC>::LaunchEvalKernel(const double* score) const {
+void CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const {
   const int num_blocks = (this->num_data_ + NUM_DATA_PER_EVAL_THREAD - 1) / NUM_DATA_PER_EVAL_THREAD;
   if (this->cuda_weights_ != nullptr) {
     EvalKernel<CUDA_METRIC, true><<<num_blocks, NUM_DATA_PER_EVAL_THREAD>>>(
@@ -50,18 +52,17 @@ double CUDARegressionMetricInterface<HOST_METRIC, CUDA_METRIC>::LaunchEvalKernel
       this->num_data_, this->cuda_labels_, this->cuda_weights_, score, reduce_block_buffer_.RawData());
   }
   ShuffleReduceSumGlobal<double, double>(reduce_block_buffer_.RawData(), num_blocks, reduce_block_buffer_inner_.RawData());
-  double sum_loss = 0.0;
-  CopyFromCUDADeviceToHost<double>(&sum_loss, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__);
-  double sum_weight = static_cast<double>(this->num_data_);
+  CopyFromCUDADeviceToHost<double>(sum_loss, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__);
+  *sum_weight = static_cast<double>(this->num_data_);
   if (this->cuda_weights_ != nullptr) {
     ShuffleReduceSumGlobal<double, double>(reduce_block_buffer_.RawData() + num_blocks, num_blocks, reduce_block_buffer_inner_.RawData());
-    CopyFromCUDADeviceToHost<double>(&sum_weight, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__);
+    CopyFromCUDADeviceToHost<double>(sum_weight, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__);
   }
-  return this->AverageLoss(sum_loss, sum_weight);
 }
 
-template double CUDARegressionMetricInterface<RMSEMetric, CUDARMSEMetric>::LaunchEvalKernel(const double* score) const;
-template double CUDARegressionMetricInterface<L2Metric, CUDAL2Metric>::LaunchEvalKernel(const double* score) const;
+template void CUDAPointwiseMetricInterface<RMSEMetric, CUDARMSEMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<L2Metric, CUDAL2Metric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<BinaryLoglossMetric, CUDABinaryLoglossMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
 
 }  // namespace LightGBM
 

src/metric/cuda/cuda_pointwise_metric.hpp

+/*!
+ * Copyright (c) 2022 Microsoft Corporation. All rights reserved.
+ * Licensed under the MIT License. See LICENSE file in the project root for
+ * license information.
+ */
+
+#ifndef LIGHTGBM_METRIC_CUDA_CUDA_POINTWISE_METRIC_HPP_
+#define LIGHTGBM_METRIC_CUDA_CUDA_POINTWISE_METRIC_HPP_
+
+#ifdef USE_CUDA_EXP
+
+#include <LightGBM/cuda/cuda_metric.hpp>
+#include <LightGBM/cuda/cuda_utils.h>
+
+#include <vector>
+
+#define NUM_DATA_PER_EVAL_THREAD (1024)
+
+namespace LightGBM {
+
+template <typename HOST_METRIC, typename CUDA_METRIC>
+class CUDAPointwiseMetricInterface: public CUDAMetricInterface<HOST_METRIC> {
+ public:
+  explicit CUDAPointwiseMetricInterface(const Config& config): CUDAMetricInterface<HOST_METRIC>(config), num_class_(config.num_class) {}
+
+  virtual ~CUDAPointwiseMetricInterface() {}
+
+  void Init(const Metadata& metadata, data_size_t num_data) override;
+
+ protected:
+  void LaunchEvalKernel(const double* score_convert, double* sum_loss, double* sum_weight) const;
+
+  mutable CUDAVector<double> score_convert_buffer_;
+  CUDAVector<double> reduce_block_buffer_;
+  CUDAVector<double> reduce_block_buffer_inner_;
+  const int num_class_;
+};
+
+}  // namespace LightGBM
+
+#endif  // USE_CUDA_EXP
+
+#endif  // LIGHTGBM_METRIC_CUDA_CUDA_POINTWISE_METRIC_HPP_

src/metric/cuda/cuda_regression_metric.cpp

@@ -12,31 +12,16 @@
 
 namespace LightGBM {
 
-template <typename HOST_METRIC, typename CUDA_METRIC>
-void CUDARegressionMetricInterface<HOST_METRIC, CUDA_METRIC>::Init(const Metadata& metadata, data_size_t num_data) {
-  CUDAMetricInterface<HOST_METRIC>::Init(metadata, num_data);
-  const int max_num_reduce_blocks = (this->num_data_ + NUM_DATA_PER_EVAL_THREAD - 1) / NUM_DATA_PER_EVAL_THREAD;
-  if (this->cuda_weights_ == nullptr) {
-    reduce_block_buffer_.Resize(max_num_reduce_blocks);
-  } else {
-    reduce_block_buffer_.Resize(max_num_reduce_blocks * 2);
-  }
-  const int max_num_reduce_blocks_inner = (max_num_reduce_blocks + NUM_DATA_PER_EVAL_THREAD - 1) / NUM_DATA_PER_EVAL_THREAD;
-  if (this->cuda_weights_ == nullptr) {
-    reduce_block_buffer_inner_.Resize(max_num_reduce_blocks_inner);
-  } else {
-    reduce_block_buffer_inner_.Resize(max_num_reduce_blocks_inner * 2);
-  }
-}
-
 template <typename HOST_METRIC, typename CUDA_METRIC>
 std::vector<double> CUDARegressionMetricInterface<HOST_METRIC, CUDA_METRIC>::Eval(const double* score, const ObjectiveFunction* objective) const {
   const double* score_convert = score;
   if (objective != nullptr && objective->NeedConvertOutputCUDA()) {
-    score_convert_buffer_.Resize(static_cast<size_t>(this->num_data_) * static_cast<size_t>(this->num_class_));
-    score_convert = objective->ConvertOutputCUDA(this->num_data_, score, score_convert_buffer_.RawData());
+    this->score_convert_buffer_.Resize(static_cast<size_t>(this->num_data_) * static_cast<size_t>(this->num_class_));
+    score_convert = objective->ConvertOutputCUDA(this->num_data_, score, this->score_convert_buffer_.RawData());
   }
-  const double eval_score = LaunchEvalKernel(score_convert);
+  double sum_loss = 0.0, sum_weight = 0.0;
+  this->LaunchEvalKernel(score_convert, &sum_loss, &sum_weight);
+  const double eval_score = this->AverageLoss(sum_loss, sum_weight);
   return std::vector<double>{eval_score};
 }
 

src/metric/cuda/cuda_regression_metric.hpp

@@ -14,30 +14,20 @@
 
 #include <vector>
 
+#include "cuda_pointwise_metric.hpp"
 #include "../regression_metric.hpp"
 
-#define NUM_DATA_PER_EVAL_THREAD (1024)
-
 namespace LightGBM {
 
 template <typename HOST_METRIC, typename CUDA_METRIC>
-class CUDARegressionMetricInterface: public CUDAMetricInterface<HOST_METRIC> {
+class CUDARegressionMetricInterface: public CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC> {
  public:
-  explicit CUDARegressionMetricInterface(const Config& config): CUDAMetricInterface<HOST_METRIC>(config), num_class_(config.num_class) {}
+  explicit CUDARegressionMetricInterface(const Config& config):
+    CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>(config) {}
 
   virtual ~CUDARegressionMetricInterface() {}
 
-  void Init(const Metadata& metadata, data_size_t num_data) override;
-
   std::vector<double> Eval(const double* score, const ObjectiveFunction* objective) const override;
-
- protected:
-  double LaunchEvalKernel(const double* score_convert) const;
-
-  mutable CUDAVector<double> score_convert_buffer_;
-  CUDAVector<double> reduce_block_buffer_;
-  CUDAVector<double> reduce_block_buffer_inner_;
-  const int num_class_;
 };
 
 class CUDARMSEMetric: public CUDARegressionMetricInterface<RMSEMetric, CUDARMSEMetric> {

src/metric/metric.cpp

@@ -11,13 +11,14 @@
 #include "regression_metric.hpp"
 #include "xentropy_metric.hpp"
 
+#include "cuda/cuda_binary_metric.hpp"
 #include "cuda/cuda_regression_metric.hpp"
 
 namespace LightGBM {
 
 Metric* Metric::CreateMetric(const std::string& type, const Config& config) {
   #ifdef USE_CUDA_EXP
-  if (config.device_type == std::string("cuda_exp")) {
+  if (config.device_type == std::string("cuda_exp") && config.boosting == std::string("gbdt")) {
     if (type == std::string("l2")) {
       return new CUDAL2Metric(config);
     } else if (type == std::string("rmse")) {
@@ -38,8 +39,7 @@ Metric* Metric::CreateMetric(const std::string& type, const Config& config) {
       Log::Warning("Metric poisson is not implemented in cuda_exp version. Fall back to evaluation on CPU.");
       return new PoissonMetric(config);
     } else if (type == std::string("binary_logloss")) {
-      Log::Warning("Metric binary_logloss is not implemented in cuda_exp version. Fall back to evaluation on CPU.");
-      return new BinaryLoglossMetric(config);
+      return new CUDABinaryLoglossMetric(config);
     } else if (type == std::string("binary_error")) {
       Log::Warning("Metric binary_error is not implemented in cuda_exp version. Fall back to evaluation on CPU.");
       return new BinaryErrorMetric(config);

src/objective/cuda/cuda_binary_objective.hpp

@@ -33,6 +33,8 @@ class CUDABinaryLogloss : public CUDAObjectiveInterface<BinaryLogloss> {
 
   void Init(const Metadata& metadata, data_size_t num_data) override;
 
+  bool NeedConvertOutputCUDA() const override { return true; }
+
  private:
   void LaunchGetGradientsKernel(const double* scores, score_t* gradients, score_t* hessians) const override;