diff --git a/include/LightGBM/cuda/cuda_utils.h b/include/LightGBM/cuda/cuda_utils.h index 771e1561f767..46ac5a9149d7 100644 --- a/include/LightGBM/cuda/cuda_utils.h +++ b/include/LightGBM/cuda/cuda_utils.h @@ -12,6 +12,7 @@ #include #include #include +#include namespace LightGBM { @@ -177,6 +178,15 @@ class CUDAVector { size_t size_; }; +template +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 index 592c1d38ca1b..85e886998d91 100644 --- a/src/metric/cuda/cuda_pointwise_metric.cpp +++ b/src/metric/cuda/cuda_pointwise_metric.cpp @@ -33,6 +33,14 @@ template void CUDAPointwiseMetricInterface::Init(con 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); +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); +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 f72569fa4816..6b200798f4c8 100644 --- a/src/metric/cuda/cuda_pointwise_metric.cu +++ b/src/metric/cuda/cuda_pointwise_metric.cu @@ -64,6 +64,14 @@ template void CUDAPointwiseMetricInterface::LaunchEv 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; +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; +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_regression_metric.cpp b/src/metric/cuda/cuda_regression_metric.cpp index a3af2d00f894..4036b7164c32 100644 --- 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(config), alpha_(config.alpha) {} +CUDAL1Metric::CUDAL1Metric(const Config& config): CUDARegressionMetricInterface(config) {} + +CUDAHuberLossMetric::CUDAHuberLossMetric(const Config& config): CUDARegressionMetricInterface(config), alpha_(config.alpha) {} + +CUDAFairLossMetric::CUDAFairLossMetric(const Config& config): CUDARegressionMetricInterface(config) , fair_c_(config.fair_c) {} + +CUDAPoissonMetric::CUDAPoissonMetric(const Config& config): CUDARegressionMetricInterface(config) {} + +CUDAMAPEMetric::CUDAMAPEMetric(const Config& config): CUDARegressionMetricInterface(config) {} + +CUDAGammaMetric::CUDAGammaMetric(const Config& config): CUDARegressionMetricInterface(config) {} + +CUDAGammaDevianceMetric::CUDAGammaDevianceMetric(const Config& config): CUDARegressionMetricInterface(config) {} + +CUDATweedieMetric::CUDATweedieMetric(const Config& config): CUDARegressionMetricInterface(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 index ab9f4f4028bc..4cfd996d837d 100644 --- a/src/metric/cuda/cuda_regression_metric.hpp +++ b/src/metric/cuda/cuda_regression_metric.hpp @@ -75,6 +75,139 @@ class CUDAQuantileMetric : public CUDARegressionMetricInterface { + 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 { + 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 { + 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 { + 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 { + 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 { + 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 { + 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 { + 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 index 4c4e046cd6f4..531ef5650f4d 100644 --- 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