From 54486b4f6bc1962e2583e05c3cedf8778bce0393 Mon Sep 17 00:00:00 2001 From: shiyu1994 Date: Thu, 16 Mar 2023 18:00:30 +0800 Subject: [PATCH] [CUDA] Add quantile metric for new CUDA version (contribute to #5163) (#5665) Co-authored-by: James Lamb --- src/metric/cuda/cuda_binary_metric.hpp | 2 +- src/metric/cuda/cuda_pointwise_metric.cpp | 1 + src/metric/cuda/cuda_pointwise_metric.cu | 11 +++++---- src/metric/cuda/cuda_pointwise_metric.hpp | 2 ++ src/metric/cuda/cuda_regression_metric.cpp | 2 ++ src/metric/cuda/cuda_regression_metric.hpp | 27 ++++++++++++++++++++-- src/metric/metric.cpp | 3 +-- tests/python_package_test/test_engine.py | 6 +++++ 8 files changed, 44 insertions(+), 10 deletions(-) diff --git a/src/metric/cuda/cuda_binary_metric.hpp b/src/metric/cuda/cuda_binary_metric.hpp index 72d9edc8b40e..86dbdce986f5 100644 --- a/src/metric/cuda/cuda_binary_metric.hpp +++ b/src/metric/cuda/cuda_binary_metric.hpp @@ -35,7 +35,7 @@ class CUDABinaryLoglossMetric: public CUDABinaryMetricInterface kEpsilon) { diff --git a/src/metric/cuda/cuda_pointwise_metric.cpp b/src/metric/cuda/cuda_pointwise_metric.cpp index bf9a0cd2867c..592c1d38ca1b 100644 --- a/src/metric/cuda/cuda_pointwise_metric.cpp +++ b/src/metric/cuda/cuda_pointwise_metric.cpp @@ -31,6 +31,7 @@ void CUDAPointwiseMetricInterface::Init(const Metadata template void CUDAPointwiseMetricInterface::Init(const Metadata& metadata, data_size_t num_data); template void CUDAPointwiseMetricInterface::Init(const Metadata& metadata, data_size_t num_data); +template void CUDAPointwiseMetricInterface::Init(const Metadata& metadata, data_size_t num_data); template void CUDAPointwiseMetricInterface::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 index 2b276bdc3e01..f72569fa4816 100644 --- a/src/metric/cuda/cuda_pointwise_metric.cu +++ b/src/metric/cuda/cuda_pointwise_metric.cu @@ -16,14 +16,14 @@ namespace LightGBM { template __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(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(point_metric, shared_mem_buffer, NUM_DATA_PER_EVAL_THREAD); if (threadIdx.x == 0) { @@ -46,10 +46,10 @@ void CUDAPointwiseMetricInterface::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<<>>( - 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<<>>( - 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(reduce_block_buffer_.RawData(), num_blocks, reduce_block_buffer_inner_.RawData()); CopyFromCUDADeviceToHost(sum_loss, reduce_block_buffer_inner_.RawData(), 1, __FILE__, __LINE__); @@ -62,6 +62,7 @@ void CUDAPointwiseMetricInterface::LaunchEvalKernel(co template void CUDAPointwiseMetricInterface::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const; template void CUDAPointwiseMetricInterface::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const; +template void CUDAPointwiseMetricInterface::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const; template void CUDAPointwiseMetricInterface::LaunchEvalKernel(const double* score, double* sum_loss, double* sum_weight) const; } // namespace LightGBM diff --git a/src/metric/cuda/cuda_pointwise_metric.hpp b/src/metric/cuda/cuda_pointwise_metric.hpp index fbdff27486d3..dae1c6a7fa68 100644 --- a/src/metric/cuda/cuda_pointwise_metric.hpp +++ b/src/metric/cuda/cuda_pointwise_metric.hpp @@ -30,6 +30,8 @@ class CUDAPointwiseMetricInterface: public CUDAMetricInterface { protected: void LaunchEvalKernel(const double* score_convert, double* sum_loss, double* sum_weight) const; + virtual double GetParamFromConfig() const { return 0.0; } + mutable CUDAVector score_convert_buffer_; CUDAVector reduce_block_buffer_; CUDAVector reduce_block_buffer_inner_; diff --git a/src/metric/cuda/cuda_regression_metric.cpp b/src/metric/cuda/cuda_regression_metric.cpp index 5e61214ad1e8..a3af2d00f894 100644 --- a/src/metric/cuda/cuda_regression_metric.cpp +++ b/src/metric/cuda/cuda_regression_metric.cpp @@ -29,6 +29,8 @@ CUDARMSEMetric::CUDARMSEMetric(const Config& config): CUDARegressionMetricInterf CUDAL2Metric::CUDAL2Metric(const Config& config): CUDARegressionMetricInterface(config) {} +CUDAQuantileMetric::CUDAQuantileMetric(const Config& config): CUDARegressionMetricInterface(config), alpha_(config.alpha) {} + } // namespace LightGBM #endif // USE_CUDA diff --git a/src/metric/cuda/cuda_regression_metric.hpp b/src/metric/cuda/cuda_regression_metric.hpp index 6e9d44a6b046..ab9f4f4028bc 100644 --- a/src/metric/cuda/cuda_regression_metric.hpp +++ b/src/metric/cuda/cuda_regression_metric.hpp @@ -36,7 +36,7 @@ class CUDARMSEMetric: public CUDARegressionMetricInterface { + 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 diff --git a/src/metric/metric.cpp b/src/metric/metric.cpp index 385eea1dcecb..4c4e046cd6f4 100644 --- a/src/metric/metric.cpp +++ b/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); diff --git a/tests/python_package_test/test_engine.py b/tests/python_package_test/test_engine.py index 4691120184aa..8633521cf02b 100644 --- a/tests/python_package_test/test_engine.py +++ b/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 train_booster(params=params_obj_verbose, **kwargs): 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