[CUDA] Add more CUDA Regression Metrics (#5924)

* add l1 metric for cuda_exp * add huber/fair metric for cuda_exp * add poisson/mape/gamma/gamma_deviance/tweedie metrics for cuda_exp * fix cpplint error * fix return error
microsoft · Jun 16, 2023 · 07e3cf4 · 07e3cf4
1 parent 30942a3
commit 07e3cf4
Show file tree

Hide file tree

Showing 6 changed files with 183 additions and 16 deletions.
diff --git a/include/LightGBM/cuda/cuda_utils.h b/include/LightGBM/cuda/cuda_utils.h
@@ -12,6 +12,7 @@
 #include <stdio.h>
 #include <LightGBM/utils/log.h>
 #include <vector>
+#include <cmath>
 
 namespace LightGBM {
 
@@ -177,6 +178,15 @@ class CUDAVector {
   size_t size_;
 };
 
+template <typename T>
+static __device__ T SafeLog(T x) {
+  if (x > 0) {
+    return std::log(x);
+  } else {
+    return -INFINITY;
+  }
+}
+
 }  // namespace LightGBM
 
 #endif  // USE_CUDA

diff --git a/src/metric/cuda/cuda_pointwise_metric.cpp b/src/metric/cuda/cuda_pointwise_metric.cpp
@@ -33,6 +33,14 @@ template void CUDAPointwiseMetricInterface<RMSEMetric, CUDARMSEMetric>::Init(con
 template void CUDAPointwiseMetricInterface<L2Metric, CUDAL2Metric>::Init(const Metadata& metadata, data_size_t num_data);
 template void CUDAPointwiseMetricInterface<QuantileMetric, CUDAQuantileMetric>::Init(const Metadata& metadata, data_size_t num_data);
 template void CUDAPointwiseMetricInterface<BinaryLoglossMetric, CUDABinaryLoglossMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<L1Metric, CUDAL1Metric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<HuberLossMetric, CUDAHuberLossMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<FairLossMetric, CUDAFairLossMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<PoissonMetric, CUDAPoissonMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<MAPEMetric, CUDAMAPEMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<GammaMetric, CUDAGammaMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<GammaDevianceMetric, CUDAGammaDevianceMetric>::Init(const Metadata& metadata, data_size_t num_data);
+template void CUDAPointwiseMetricInterface<TweedieMetric, CUDATweedieMetric>::Init(const Metadata& metadata, data_size_t num_data);
 
 }  // namespace LightGBM
 

diff --git a/src/metric/cuda/cuda_pointwise_metric.cu b/src/metric/cuda/cuda_pointwise_metric.cu
@@ -64,6 +64,14 @@ template void CUDAPointwiseMetricInterface<RMSEMetric, CUDARMSEMetric>::LaunchEv
 template void CUDAPointwiseMetricInterface<L2Metric, CUDAL2Metric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
 template void CUDAPointwiseMetricInterface<QuantileMetric, CUDAQuantileMetric>::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;
+template void CUDAPointwiseMetricInterface<L1Metric, CUDAL1Metric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<HuberLossMetric, CUDAHuberLossMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<FairLossMetric, CUDAFairLossMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<PoissonMetric, CUDAPoissonMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<MAPEMetric, CUDAMAPEMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<GammaMetric, CUDAGammaMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<GammaDevianceMetric, CUDAGammaDevianceMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
+template void CUDAPointwiseMetricInterface<TweedieMetric, CUDATweedieMetric>::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const;
 
 }  // namespace LightGBM
 

diff --git a/src/metric/cuda/cuda_regression_metric.cpp b/src/metric/cuda/cuda_regression_metric.cpp
@@ -31,6 +31,22 @@ CUDAL2Metric::CUDAL2Metric(const Config& config): CUDARegressionMetricInterface<
 
 CUDAQuantileMetric::CUDAQuantileMetric(const Config& config): CUDARegressionMetricInterface<QuantileMetric, CUDAQuantileMetric>(config), alpha_(config.alpha) {}
 
+CUDAL1Metric::CUDAL1Metric(const Config& config): CUDARegressionMetricInterface<L1Metric, CUDAL1Metric>(config) {}
+
+CUDAHuberLossMetric::CUDAHuberLossMetric(const Config& config): CUDARegressionMetricInterface<HuberLossMetric, CUDAHuberLossMetric>(config), alpha_(config.alpha) {}
+
+CUDAFairLossMetric::CUDAFairLossMetric(const Config& config): CUDARegressionMetricInterface<FairLossMetric, CUDAFairLossMetric>(config) , fair_c_(config.fair_c) {}
+
+CUDAPoissonMetric::CUDAPoissonMetric(const Config& config): CUDARegressionMetricInterface<PoissonMetric, CUDAPoissonMetric>(config) {}
+
+CUDAMAPEMetric::CUDAMAPEMetric(const Config& config): CUDARegressionMetricInterface<MAPEMetric, CUDAMAPEMetric>(config) {}
+
+CUDAGammaMetric::CUDAGammaMetric(const Config& config): CUDARegressionMetricInterface<GammaMetric, CUDAGammaMetric>(config) {}
+
+CUDAGammaDevianceMetric::CUDAGammaDevianceMetric(const Config& config): CUDARegressionMetricInterface<GammaDevianceMetric, CUDAGammaDevianceMetric>(config) {}
+
+CUDATweedieMetric::CUDATweedieMetric(const Config& config): CUDARegressionMetricInterface<TweedieMetric, CUDATweedieMetric>(config) , tweedie_variance_power_(config.tweedie_variance_power) {}
+
 }  // namespace LightGBM
 
 #endif  // USE_CUDA
diff --git a/src/metric/cuda/cuda_regression_metric.hpp b/src/metric/cuda/cuda_regression_metric.hpp
@@ -75,6 +75,139 @@ class CUDAQuantileMetric : public CUDARegressionMetricInterface<QuantileMetric,
   const double alpha_;
 };
 
+class CUDAL1Metric : public CUDARegressionMetricInterface<L1Metric, CUDAL1Metric> {
+ public:
+  explicit CUDAL1Metric(const Config& config);
+
+  virtual ~CUDAL1Metric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double /*alpha*/) {
+    return std::fabs(score - label);
+  }
+};
+
+class CUDAHuberLossMetric : public CUDARegressionMetricInterface<HuberLossMetric, CUDAHuberLossMetric> {
+ public:
+  explicit CUDAHuberLossMetric(const Config& config);
+
+  virtual ~CUDAHuberLossMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double alpha) {
+    const double diff = score - label;
+    if (std::abs(diff) <= alpha) {
+      return 0.5f * diff * diff;
+    } else {
+      return alpha * (std::abs(diff) - 0.5f * alpha);
+    }
+  }
+
+  double GetParamFromConfig() const override {
+    return alpha_;
+  }
+ private:
+  const double alpha_;
+};
+
+class CUDAFairLossMetric : public CUDARegressionMetricInterface<FairLossMetric, CUDAFairLossMetric> {
+ public:
+  explicit CUDAFairLossMetric(const Config& config);
+
+  virtual ~CUDAFairLossMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double fair_c) {
+    const double x = std::fabs(score - label);
+    const double c =  fair_c;
+    return c * x - c * c * std::log1p(x / c);
+  }
+
+  double GetParamFromConfig() const override {
+    return fair_c_;
+  }
+
+ private:
+  const double fair_c_;
+};
+
+class CUDAPoissonMetric : public CUDARegressionMetricInterface<PoissonMetric, CUDAPoissonMetric> {
+ public:
+  explicit CUDAPoissonMetric(const Config& config);
+
+  virtual ~CUDAPoissonMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double /*alpha*/) {
+    const double eps = 1e-10f;
+    if (score < eps) {
+      score = eps;
+    }
+    return score - label * std::log(score);
+  }
+};
+
+class CUDAMAPEMetric : public CUDARegressionMetricInterface<MAPEMetric, CUDAMAPEMetric> {
+ public:
+  explicit CUDAMAPEMetric(const Config& config);
+
+  virtual ~CUDAMAPEMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double /*alpha*/) {
+    return std::fabs((label - score)) / fmax(1.0f, std::fabs(label));
+  }
+};
+
+class CUDAGammaMetric : public CUDARegressionMetricInterface<GammaMetric, CUDAGammaMetric> {
+ public:
+  explicit CUDAGammaMetric(const Config& config);
+
+  virtual ~CUDAGammaMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double /*alpha*/) {
+    const double psi = 1.0;
+    const double theta = -1.0 / score;
+    const double a = psi;
+    const double b = -SafeLog(-theta);
+    const double c = 1. / psi * SafeLog(label / psi) - SafeLog(label) - 0;  // 0 = std::lgamma(1.0 / psi) = std::lgamma(1.0);
+    return -((label * theta - b) / a + c);
+  }
+};
+
+class CUDAGammaDevianceMetric : public CUDARegressionMetricInterface<GammaDevianceMetric, CUDAGammaDevianceMetric> {
+ public:
+  explicit CUDAGammaDevianceMetric(const Config& config);
+
+  virtual ~CUDAGammaDevianceMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double /*alpha*/) {
+    const double epsilon = 1.0e-9;
+    const double tmp = label / (score + epsilon);
+    return tmp - SafeLog(tmp) - 1;
+  }
+};
+
+class CUDATweedieMetric : public CUDARegressionMetricInterface<TweedieMetric, CUDATweedieMetric> {
+ public:
+  explicit CUDATweedieMetric(const Config& config);
+
+  virtual ~CUDATweedieMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score,  double tweedie_variance_power) {
+    const double rho = tweedie_variance_power;
+    const double eps = 1e-10f;
+    if (score < eps) {
+      score = eps;
+    }
+    const double a = label * std::exp((1 - rho) * std::log(score)) / (1 - rho);
+    const double b = std::exp((2 - rho) * std::log(score)) / (2 - rho);
+    return -a + b;
+  }
+
+  double GetParamFromConfig() const override {
+    return tweedie_variance_power_;
+  }
+
+ private:
+  const double tweedie_variance_power_;
+};
+
 }  // namespace LightGBM
 
 #endif  // USE_CUDA

diff --git a/src/metric/metric.cpp b/src/metric/metric.cpp
@@ -24,19 +24,15 @@ Metric* Metric::CreateMetric(const std::string& type, const Config& config) {
     } else if (type == std::string("rmse")) {
       return new CUDARMSEMetric(config);
     } else if (type == std::string("l1")) {
-      Log::Warning("Metric l1 is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new L1Metric(config);
+      return new CUDAL1Metric(config);
     } else if (type == std::string("quantile")) {
       return new CUDAQuantileMetric(config);
     } else if (type == std::string("huber")) {
-      Log::Warning("Metric huber is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new HuberLossMetric(config);
+      return new CUDAHuberLossMetric(config);
     } else if (type == std::string("fair")) {
-      Log::Warning("Metric fair is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new FairLossMetric(config);
+      return new CUDAFairLossMetric(config);
     } else if (type == std::string("poisson")) {
-      Log::Warning("Metric poisson is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new PoissonMetric(config);
+      return new CUDAPoissonMetric(config);
     } else if (type == std::string("binary_logloss")) {
       return new CUDABinaryLoglossMetric(config);
     } else if (type == std::string("binary_error")) {
@@ -73,17 +69,13 @@ Metric* Metric::CreateMetric(const std::string& type, const Config& config) {
       Log::Warning("Metric kullback_leibler is not implemented in cuda version. Fall back to evaluation on CPU.");
       return new KullbackLeiblerDivergence(config);
     } else if (type == std::string("mape")) {
-      Log::Warning("Metric mape is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new MAPEMetric(config);
+      return new CUDAMAPEMetric(config);
     } else if (type == std::string("gamma")) {
-      Log::Warning("Metric gamma is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new GammaMetric(config);
+      return new CUDAGammaMetric(config);
     } else if (type == std::string("gamma_deviance")) {
-      Log::Warning("Metric gamma_deviance is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new GammaDevianceMetric(config);
+      return new CUDAGammaDevianceMetric(config);
     } else if (type == std::string("tweedie")) {
-      Log::Warning("Metric tweedie is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new TweedieMetric(config);
+      return new CUDATweedieMetric(config);
     }
   } else {
   #endif  // USE_CUDA