[CUDA] Add quantile metric for new CUDA version (contribute to #5163) (#5665)

Co-authored-by: James Lamb <jaylamb20@gmail.com>

src/metric/cuda/cuda_binary_metric.hpp

@@ -35,7 +35,7 @@ class CUDABinaryLoglossMetric: public CUDABinaryMetricInterface<BinaryLoglossMet
 
   virtual ~CUDABinaryLoglossMetric() {}
 
-  __device__ static double MetricOnPointCUDA(label_t label, double score) {
+  __device__ static double MetricOnPointCUDA(label_t label, double score, const double /*param*/) {
     // score should have been converted to probability
     if (label <= 0) {
       if (1.0f - score > kEpsilon) {

src/metric/cuda/cuda_pointwise_metric.cpp

@@ -31,6 +31,7 @@ void CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>::Init(const Metadata
 
 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<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);
 
 }  // namespace LightGBM

src/metric/cuda/cuda_pointwise_metric.cu

@@ -16,14 +16,14 @@ namespace LightGBM {
 
 template <typename CUDA_METRIC, bool USE_WEIGHTS>
 __global__ void EvalKernel(const data_size_t num_data, const label_t* labels, const label_t* weights,
-                           const double* scores, double* reduce_block_buffer) {
+                           const double* scores, double* reduce_block_buffer, const double param) {
   __shared__ double shared_mem_buffer[32];
   const data_size_t index = static_cast<data_size_t>(threadIdx.x + blockIdx.x * blockDim.x);
   double point_metric = 0.0;
   if (index < num_data) {
     point_metric = USE_WEIGHTS ?
-      CUDA_METRIC::MetricOnPointCUDA(labels[index], scores[index]) * weights[index] :
-      CUDA_METRIC::MetricOnPointCUDA(labels[index], scores[index]);
+      CUDA_METRIC::MetricOnPointCUDA(labels[index], scores[index], param) * weights[index] :
+      CUDA_METRIC::MetricOnPointCUDA(labels[index], scores[index], param);
   }
   const double block_sum_point_metric = ShuffleReduceSum<double>(point_metric, shared_mem_buffer, NUM_DATA_PER_EVAL_THREAD);
   if (threadIdx.x == 0) {
@@ -46,10 +46,10 @@ void CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>::LaunchEvalKernel(co
   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>>>(
-      this->num_data_, this->cuda_labels_, this->cuda_weights_, score, reduce_block_buffer_.RawData());
+      this->num_data_, this->cuda_labels_, this->cuda_weights_, score, reduce_block_buffer_.RawData(), GetParamFromConfig());
   } else {
     EvalKernel<CUDA_METRIC, false><<<num_blocks, NUM_DATA_PER_EVAL_THREAD>>>(
-      this->num_data_, this->cuda_labels_, this->cuda_weights_, score, reduce_block_buffer_.RawData());
+      this->num_data_, this->cuda_labels_, this->cuda_weights_, score, reduce_block_buffer_.RawData(), GetParamFromConfig());
   }
   ShuffleReduceSumGlobal<double, double>(reduce_block_buffer_.RawData(), num_blocks, reduce_block_buffer_inner_.RawData());
   CopyFromCUDADeviceToHost<double>(sum_loss, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__);
@@ -62,6 +62,7 @@ void CUDAPointwiseMetricInterface<HOST_METRIC, CUDA_METRIC>::LaunchEvalKernel(co
 
 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<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;
 
 }  // namespace LightGBM

src/metric/cuda/cuda_pointwise_metric.hpp

@@ -30,6 +30,8 @@ class CUDAPointwiseMetricInterface: public CUDAMetricInterface<HOST_METRIC> {
  protected:
   void LaunchEvalKernel(const double* score_convert, double* sum_loss, double* sum_weight) const;
 
+  virtual double GetParamFromConfig() const { return 0.0; }
+
   mutable CUDAVector<double> score_convert_buffer_;
   CUDAVector<double> reduce_block_buffer_;
   CUDAVector<double> reduce_block_buffer_inner_;

src/metric/cuda/cuda_regression_metric.cpp

@@ -29,6 +29,8 @@ CUDARMSEMetric::CUDARMSEMetric(const Config& config): CUDARegressionMetricInterf
 
 CUDAL2Metric::CUDAL2Metric(const Config& config): CUDARegressionMetricInterface<L2Metric, CUDAL2Metric>(config) {}
 
+CUDAQuantileMetric::CUDAQuantileMetric(const Config& config): CUDARegressionMetricInterface<QuantileMetric, CUDAQuantileMetric>(config), alpha_(config.alpha) {}
+
 }  // namespace LightGBM
 
 #endif  // USE_CUDA

src/metric/cuda/cuda_regression_metric.hpp

@@ -36,7 +36,7 @@ class CUDARMSEMetric: public CUDARegressionMetricInterface<RMSEMetric, CUDARMSEM
 
   virtual ~CUDARMSEMetric() {}
 
-  __device__ inline static double MetricOnPointCUDA(label_t label, double score) {
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score, double /*alpha*/) {
     return (score - label) * (score - label);
   }
 };
@@ -47,11 +47,34 @@ class CUDAL2Metric : public CUDARegressionMetricInterface<L2Metric, CUDAL2Metric
 
   virtual ~CUDAL2Metric() {}
 
-  __device__ inline static double MetricOnPointCUDA(label_t label, double score) {
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score, double /*alpha*/) {
     return (score - label) * (score - label);
   }
 };
 
+class CUDAQuantileMetric : public CUDARegressionMetricInterface<QuantileMetric, CUDAQuantileMetric> {
+ public:
+  explicit CUDAQuantileMetric(const Config& config);
+
+  virtual ~CUDAQuantileMetric() {}
+
+  __device__ inline static double MetricOnPointCUDA(label_t label, double score, double alpha) {
+    double delta = label - score;
+    if (delta < 0) {
+      return (alpha - 1.0f) * delta;
+    } else {
+      return alpha * delta;
+    }
+  }
+
+  double GetParamFromConfig() const override {
+    return alpha_;
+  }
+
+ private:
+  const double alpha_;
+};
+
 }  // namespace LightGBM
 
 #endif  // USE_CUDA

src/metric/metric.cpp

@@ -27,8 +27,7 @@ Metric* Metric::CreateMetric(const std::string& type, const Config& config) {
       Log::Warning("Metric l1 is not implemented in cuda version. Fall back to evaluation on CPU.");
       return new L1Metric(config);
     } else if (type == std::string("quantile")) {
-      Log::Warning("Metric quantile is not implemented in cuda version. Fall back to evaluation on CPU.");
-      return new QuantileMetric(config);
+      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);

tests/python_package_test/test_engine.py

@@ -2031,6 +2031,7 @@ def test_metrics():
     params_obj_metric_log_verbose = {'objective': 'binary', 'metric': 'binary_logloss', 'verbose': -1}
     params_obj_metric_err_verbose = {'objective': 'binary', 'metric': 'binary_error', 'verbose': -1}
     params_obj_metric_inv_verbose = {'objective': 'binary', 'metric': 'invalid_metric', 'verbose': -1}
+    params_obj_metric_quant_verbose = {'objective': 'regression', 'metric': 'quantile', 'verbose': 2}
     params_obj_metric_multi_verbose = {'objective': 'binary',
                                        'metric': ['binary_logloss', 'binary_error'],
                                        'verbose': -1}
@@ -2080,6 +2081,11 @@ def test_metrics():
     assert len(res) == 2
     assert 'valid binary_error-mean' in res
 
+    # metric in args overwrites one in params
+    res = get_cv_result(params=params_obj_metric_quant_verbose)
+    assert len(res) == 2
+    assert 'valid quantile-mean' in res
+
     # multiple metrics in params
     res = get_cv_result(params=params_obj_metric_multi_verbose)
     assert len(res) == 4