From 8dd0debb94c42b134a269ce27e2eca5c62a83744 Mon Sep 17 00:00:00 2001 From: Max Buckley Date: Tue, 2 Jun 2026 00:38:19 +0200 Subject: [PATCH] [cuda] implement basic monotone constraint enforcement on the CUDA tree learner Implements the "basic" monotone constraint method on CUDA so that device_type="cuda" enforces monotone_constraints like device_type="cpu": - CalculateSplittedLeafOutputMC / GetSplitGainsMC device functions mirror the CPU USE_MC formulas: candidate child outputs are clamped into the leaf's [min, max] constraint range and splits whose clamped outputs violate the requested monotone direction are rejected (gain 0) - each SplitFindTask carries its feature's monotone type (real-feature-indexed config mapped to inner feature indices) - per-leaf [min, max] constraint ranges are tracked on host, mirroring BasicLeafConstraints::Update midpoint propagation, and passed to the split kernels for the (smaller, larger) leaf pair each iteration - the leaf VALUES stored in CUDASplitInfo are the clamped outputs (used for the tree and for constraint propagation), while the leaf GAINS keep their unconstrained values (they form the child's parent_gain baseline, matching the CPU BeforeNumerical gain_shift computation) - unsupported monotone configurations on CUDA fail fast with a clear error: monotone_constraints_method=intermediate/advanced, monotone_penalty>0, and use_quantized_grad+monotone This replaces the previous Log::Fatal guard: basic-method monotone constraints now work on CUDA instead of being rejected. Before this change CUDA silently ignored monotone constraints: on a 600x3 regression with constraints [1,1,1] over 100 rounds, the CUDA model violated monotonicity at 456 of 4500 grid points checked (worst violation 0.39), with predictions diverging from CPU by up to 2.8. After this change CUDA models have ZERO monotonicity violations across all tested constraint vectors, round counts and seeds, with training MSE matching CPU within 3%. Note on tree-structure parity: CPU and CUDA can build different (both valid) constrained trees because CPU additionally prunes candidate features through its FeatureHistogram::is_splittable_ cache, a search-space heuristic whose exact replication on GPU would require per-(leaf, feature) state tracking with additional host-device round trips. The constraint guarantee itself - the property users rely on - is enforced identically, predictions match bit-for-bit when constraints are inactive ([0,0,0]), and constrained training quality is equivalent (MSE within 3% of CPU, sometimes better). Adds 18 regression tests to test_dual.py: 12 enforcement cases (zero violations required), 3 CPU-quality-equivalence cases, 2 noop-constraint bit-parity cases, and 1 unsupported-config rejection case. Co-Authored-By: Claude Opus 4.8 (1M context) (cherry picked from commit 9349df83764d50b56b6508419fa583dbd7cd6413) --- src/io/config.cpp | 15 + .../cuda/cuda_best_split_finder.cpp | 31 +- .../cuda/cuda_best_split_finder.cu | 294 ++++++++++++++---- .../cuda/cuda_best_split_finder.hpp | 18 +- src/treelearner/cuda/cuda_leaf_splits.hpp | 54 ++++ .../cuda/cuda_single_gpu_tree_learner.cpp | 97 +++++- .../cuda/cuda_single_gpu_tree_learner.hpp | 16 + tests/python_package_test/test_dual.py | 126 ++++++++ 8 files changed, 583 insertions(+), 68 deletions(-) diff --git a/src/io/config.cpp b/src/io/config.cpp index 7f5b0eba46c0..59ebc042c4cf 100644 --- a/src/io/config.cpp +++ b/src/io/config.cpp @@ -461,6 +461,21 @@ void Config::CheckParamConflict(const std::unordered_map 0 && monotone_penalty >= max_depth) { Log::Warning("Monotone penalty greater than tree depth. Monotone features won't be used."); } + if (device_type == std::string("cuda") && !monotone_constraints.empty()) { + if (monotone_constraints_method != std::string("basic")) { + Log::Fatal("CUDA only supports the \"basic\" monotone_constraints_method. " + "Got \"%s\". Use device_type=cpu for intermediate/advanced methods.", + monotone_constraints_method.c_str()); + } + if (monotone_penalty > 0.0) { + Log::Fatal("monotone_penalty is not supported with device_type=cuda. " + "Set monotone_penalty=0 or use device_type=cpu."); + } + if (use_quantized_grad) { + Log::Fatal("monotone_constraints is not supported with use_quantized_grad on " + "device_type=cuda. Disable one of them or use device_type=cpu."); + } + } if (min_data_in_leaf <= 0 && min_sum_hessian_in_leaf <= kEpsilon) { Log::Warning( "Cannot set both min_data_in_leaf and min_sum_hessian_in_leaf to 0. " diff --git a/src/treelearner/cuda/cuda_best_split_finder.cpp b/src/treelearner/cuda/cuda_best_split_finder.cpp index 3e4e6203e599..821ca675d38b 100644 --- a/src/treelearner/cuda/cuda_best_split_finder.cpp +++ b/src/treelearner/cuda/cuda_best_split_finder.cpp @@ -44,6 +44,18 @@ CUDABestSplitFinder::CUDABestSplitFinder( if (has_categorical_feature_ && config->use_quantized_grad) { Log::Fatal("Quantized training on GPU with categorical features is not supported yet."); } + // Build a per-inner-feature monotone constraint vector. config->monotone_constraints + // is indexed by REAL feature index (see FeatureHistogram feature-meta init), so map + // each inner feature through RealFeatureIndex. + use_monotone_constraints_ = !config->monotone_constraints.empty(); + monotone_constraints_.resize(num_features_, 0); + if (use_monotone_constraints_) { + for (int inner_feature_index = 0; inner_feature_index < num_features_; ++inner_feature_index) { + const int real_feature_index = train_data->RealFeatureIndex(inner_feature_index); + monotone_constraints_[inner_feature_index] = + config->monotone_constraints[real_feature_index]; + } + } } CUDABestSplitFinder::~CUDABestSplitFinder() { @@ -147,6 +159,8 @@ void CUDABestSplitFinder::InitCUDAFeatureMetaInfo() { new_task->mfb_offset = feature_mfb_offsets_[inner_feature_index]; new_task->default_bin = feature_default_bins_[inner_feature_index]; new_task->num_bin = num_bin; + new_task->monotone_type = use_monotone_constraints_ ? + monotone_constraints_[inner_feature_index] : 0; ++cur_task_index; new_task = &split_find_tasks_[cur_task_index]; @@ -162,6 +176,8 @@ void CUDABestSplitFinder::InitCUDAFeatureMetaInfo() { new_task->default_bin = feature_default_bins_[inner_feature_index]; new_task->mfb_offset = feature_mfb_offsets_[inner_feature_index]; new_task->num_bin = num_bin; + new_task->monotone_type = use_monotone_constraints_ ? + monotone_constraints_[inner_feature_index] : 0; ++cur_task_index; } else { SplitFindTask* new_task = &split_find_tasks_[cur_task_index]; @@ -177,6 +193,8 @@ void CUDABestSplitFinder::InitCUDAFeatureMetaInfo() { new_task->mfb_offset = feature_mfb_offsets_[inner_feature_index]; new_task->default_bin = feature_default_bins_[inner_feature_index]; new_task->num_bin = num_bin; + new_task->monotone_type = use_monotone_constraints_ ? + monotone_constraints_[inner_feature_index] : 0; ++cur_task_index; new_task = &split_find_tasks_[cur_task_index]; @@ -192,6 +210,8 @@ void CUDABestSplitFinder::InitCUDAFeatureMetaInfo() { new_task->mfb_offset = feature_mfb_offsets_[inner_feature_index]; new_task->default_bin = feature_default_bins_[inner_feature_index]; new_task->num_bin = num_bin; + new_task->monotone_type = use_monotone_constraints_ ? + monotone_constraints_[inner_feature_index] : 0; ++cur_task_index; } } else { @@ -218,6 +238,10 @@ void CUDABestSplitFinder::InitCUDAFeatureMetaInfo() { new_task.mfb_offset = feature_mfb_offsets_[inner_feature_index]; new_task.default_bin = feature_default_bins_[inner_feature_index]; new_task.num_bin = num_bin; + // categorical features carry no monotone semantics (matching the CPU path, + // where monotone constraints only affect numerical splits) + new_task.monotone_type = (use_monotone_constraints_ && !new_task.is_categorical) ? + monotone_constraints_[inner_feature_index] : 0; ++cur_task_index; } } @@ -346,6 +370,10 @@ void CUDABestSplitFinder::FindBestSplitsForLeaf( const uint8_t larger_num_bits_in_histogram_bins, const bool smaller_leaf_below_max_depth, const bool larger_leaf_below_max_depth, + const double smaller_leaf_constraint_min, + const double smaller_leaf_constraint_max, + const double larger_leaf_constraint_min, + const double larger_leaf_constraint_max, const bool synchronize) { const bool is_smaller_leaf_valid = (num_data_in_smaller_leaf > min_data_in_leaf_ && sum_hessians_in_smaller_leaf > min_sum_hessian_in_leaf_ && @@ -359,7 +387,8 @@ void CUDABestSplitFinder::FindBestSplitsForLeaf( grad_scale, hess_scale, smaller_num_bits_in_histogram_bins, larger_num_bits_in_histogram_bins, num_data_in_smaller_leaf, num_data_in_larger_leaf); } else { LaunchFindBestSplitsForLeafKernel(smaller_leaf_splits, larger_leaf_splits, - smaller_leaf_index, larger_leaf_index, is_smaller_leaf_valid, is_larger_leaf_valid, num_data_in_smaller_leaf, num_data_in_larger_leaf); + smaller_leaf_index, larger_leaf_index, is_smaller_leaf_valid, is_larger_leaf_valid, num_data_in_smaller_leaf, num_data_in_larger_leaf, + smaller_leaf_constraint_min, smaller_leaf_constraint_max, larger_leaf_constraint_min, larger_leaf_constraint_max); } global_timer.Start("CUDABestSplitFinder::LaunchSyncBestSplitForLeafKernel"); LaunchSyncBestSplitForLeafKernel(smaller_leaf_index, larger_leaf_index, is_smaller_leaf_valid, is_larger_leaf_valid); diff --git a/src/treelearner/cuda/cuda_best_split_finder.cu b/src/treelearner/cuda/cuda_best_split_finder.cu index 24b03fb5df37..0dd3a4e9b393 100644 --- a/src/treelearner/cuda/cuda_best_split_finder.cu +++ b/src/treelearner/cuda/cuda_best_split_finder.cu @@ -12,6 +12,7 @@ #include #include +#include #include #include @@ -139,7 +140,7 @@ __device__ int ReduceBestGainForLeaves(double gain, int leaf_index, double* shar return leaf_index; } -template +template __device__ void FindBestSplitsForLeafKernelInner( // input feature information const hist_t* feature_hist_ptr, @@ -159,8 +160,12 @@ __device__ void FindBestSplitsForLeafKernelInner( const double sum_hessians, const data_size_t num_data, const double parent_output, + // monotone constraint information for this leaf + const double leaf_constraint_min, + const double leaf_constraint_max, // output parameters CUDASplitInfo* cuda_best_split_info) { + const int8_t monotone_constraint = task->monotone_type; const double cnt_factor = num_data / sum_hessians; const double min_gain_shift = parent_gain + min_gain_to_split; @@ -237,10 +242,16 @@ __device__ void FindBestSplitsForLeafKernelInner( if (sum_left_hessian >= min_sum_hessian_in_leaf && left_count >= min_data_in_leaf && sum_right_hessian >= min_sum_hessian_in_leaf && right_count >= min_data_in_leaf && (!USE_RAND || static_cast(task->num_bin - 2 - threadIdx_x) == rand_threshold)) { - double current_gain = CUDALeafSplits::GetSplitGains( - sum_left_gradient, sum_left_hessian, sum_right_gradient, - sum_right_hessian, lambda_l1, - lambda_l2, path_smooth, left_count, right_count, parent_output); + double current_gain = USE_MC ? + CUDALeafSplits::GetSplitGainsMC( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output, + leaf_constraint_min, leaf_constraint_max, monotone_constraint) : + CUDALeafSplits::GetSplitGains( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output); // gain with split is worse than without split if (current_gain > min_gain_shift) { local_gain = current_gain - min_gain_shift; @@ -261,10 +272,16 @@ __device__ void FindBestSplitsForLeafKernelInner( if (sum_left_hessian >= min_sum_hessian_in_leaf && left_count >= min_data_in_leaf && sum_right_hessian >= min_sum_hessian_in_leaf && right_count >= min_data_in_leaf && (!USE_RAND || static_cast(threadIdx_x + task->mfb_offset) == rand_threshold)) { - double current_gain = CUDALeafSplits::GetSplitGains( - sum_left_gradient, sum_left_hessian, sum_right_gradient, - sum_right_hessian, lambda_l1, - lambda_l2, path_smooth, left_count, right_count, parent_output); + double current_gain = USE_MC ? + CUDALeafSplits::GetSplitGainsMC( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output, + leaf_constraint_min, leaf_constraint_max, monotone_constraint) : + CUDALeafSplits::GetSplitGains( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output); // gain with split is worse than without split if (current_gain > min_gain_shift) { local_gain = current_gain - min_gain_shift; @@ -294,10 +311,25 @@ __device__ void FindBestSplitsForLeafKernelInner( const double sum_left_gradient = sum_gradients - sum_right_gradient; const double sum_left_hessian = sum_hessians - sum_right_hessian - kEpsilon; const data_size_t left_count = num_data - right_count; - const double left_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); - const double right_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + // Unconstrained outputs first. The leaf VALUE is clamped into the constraint + // range (MC), but the leaf GAIN stored for future splits must stay + // unconstrained: it becomes the child's parent_gain / min_gain_shift baseline, + // and the CPU reference (BeforeNumerical) recomputes that baseline as the + // unconstrained GetLeafGain of the leaf's sums. + const double left_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, + sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); + const double right_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, + sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + const double left_output = USE_MC ? + (left_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (left_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : left_output_unconstrained)) : + left_output_unconstrained; + const double right_output = USE_MC ? + (right_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (right_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : right_output_unconstrained)) : + right_output_unconstrained; cuda_best_split_info->left_sum_gradients = sum_left_gradient; cuda_best_split_info->left_sum_hessians = sum_left_hessian; cuda_best_split_info->left_count = left_count; @@ -306,10 +338,10 @@ __device__ void FindBestSplitsForLeafKernelInner( cuda_best_split_info->right_count = right_count; cuda_best_split_info->left_value = left_output; cuda_best_split_info->left_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, left_output); + sum_left_hessian, lambda_l1, lambda_l2, left_output_unconstrained); cuda_best_split_info->right_value = right_output; cuda_best_split_info->right_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, right_output); + sum_right_hessian, lambda_l1, lambda_l2, right_output_unconstrained); } else { const double sum_left_gradient = local_grad_hist; const double sum_left_hessian = local_hess_hist - kEpsilon; @@ -317,10 +349,25 @@ __device__ void FindBestSplitsForLeafKernelInner( const double sum_right_gradient = sum_gradients - sum_left_gradient; const double sum_right_hessian = sum_hessians - sum_left_hessian - kEpsilon; const data_size_t right_count = num_data - left_count; - const double left_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); - const double right_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + // Unconstrained outputs first. The leaf VALUE is clamped into the constraint + // range (MC), but the leaf GAIN stored for future splits must stay + // unconstrained: it becomes the child's parent_gain / min_gain_shift baseline, + // and the CPU reference (BeforeNumerical) recomputes that baseline as the + // unconstrained GetLeafGain of the leaf's sums. + const double left_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, + sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); + const double right_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, + sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + const double left_output = USE_MC ? + (left_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (left_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : left_output_unconstrained)) : + left_output_unconstrained; + const double right_output = USE_MC ? + (right_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (right_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : right_output_unconstrained)) : + right_output_unconstrained; cuda_best_split_info->left_sum_gradients = sum_left_gradient; cuda_best_split_info->left_sum_hessians = sum_left_hessian; cuda_best_split_info->left_count = left_count; @@ -329,10 +376,10 @@ __device__ void FindBestSplitsForLeafKernelInner( cuda_best_split_info->right_count = right_count; cuda_best_split_info->left_value = left_output; cuda_best_split_info->left_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, left_output); + sum_left_hessian, lambda_l1, lambda_l2, left_output_unconstrained); cuda_best_split_info->right_value = right_output; cuda_best_split_info->right_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, right_output); + sum_right_hessian, lambda_l1, lambda_l2, right_output_unconstrained); } } } @@ -787,7 +834,7 @@ __device__ void FindBestSplitsForLeafKernelCategoricalInner( } } -template +template __global__ void FindBestSplitsForLeafKernel( // input feature information const int8_t* is_feature_used_bytree, @@ -813,7 +860,14 @@ __global__ void FindBestSplitsForLeafKernel( CUDASplitInfo* cuda_best_split_info, // global num data in leaf const data_size_t global_num_data_in_smaller_leaf, - const data_size_t global_num_data_in_larger_leaf) { + const data_size_t global_num_data_in_larger_leaf, + // monotone leaf constraints + const double smaller_leaf_constraint_min, + const double smaller_leaf_constraint_max, + const double larger_leaf_constraint_min, + const double larger_leaf_constraint_max) { + const double leaf_constraint_min = IS_LARGER ? larger_leaf_constraint_min : smaller_leaf_constraint_min; + const double leaf_constraint_max = IS_LARGER ? larger_leaf_constraint_max : smaller_leaf_constraint_max; const unsigned int task_index = blockIdx.x; const SplitFindTask* task = tasks + task_index; const int inner_feature_index = task->inner_feature_index; @@ -856,7 +910,7 @@ __global__ void FindBestSplitsForLeafKernel( out); } else { if (!task->reverse) { - FindBestSplitsForLeafKernelInner( + FindBestSplitsForLeafKernelInner( // input feature information hist_ptr, // input task information @@ -875,10 +929,13 @@ __global__ void FindBestSplitsForLeafKernel( sum_hessians, num_data, parent_output, + // monotone constraint information + leaf_constraint_min, + leaf_constraint_max, // output parameters out); } else { - FindBestSplitsForLeafKernelInner( + FindBestSplitsForLeafKernelInner( // input feature information hist_ptr, // input task information @@ -897,6 +954,9 @@ __global__ void FindBestSplitsForLeafKernel( sum_hessians, num_data, parent_output, + // monotone constraint information + leaf_constraint_min, + leaf_constraint_max, // output parameters out); } @@ -1081,7 +1141,7 @@ __global__ void FindBestSplitsDiscretizedForLeafKernel( } } -template +template __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( // input feature information const hist_t* feature_hist_ptr, @@ -1101,11 +1161,15 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( const double sum_hessians, const data_size_t num_data, const double parent_output, + // monotone constraint information for this leaf + const double leaf_constraint_min, + const double leaf_constraint_max, // output parameters CUDASplitInfo* cuda_best_split_info, // buffer hist_t* hist_grad_buffer_ptr, hist_t* hist_hess_buffer_ptr) { + const int8_t monotone_constraint = task->monotone_type; const double cnt_factor = num_data / sum_hessians; const double min_gain_shift = parent_gain + min_gain_to_split; @@ -1197,10 +1261,16 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( if (sum_left_hessian >= min_sum_hessian_in_leaf && left_count >= min_data_in_leaf && sum_right_hessian >= min_sum_hessian_in_leaf && right_count >= min_data_in_leaf && (!USE_RAND || static_cast(task->num_bin - 2 - bin) == rand_threshold)) { - double current_gain = CUDALeafSplits::GetSplitGains( - sum_left_gradient, sum_left_hessian, sum_right_gradient, - sum_right_hessian, lambda_l1, - lambda_l2, path_smooth, left_count, right_count, parent_output); + double current_gain = USE_MC ? + CUDALeafSplits::GetSplitGainsMC( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output, + leaf_constraint_min, leaf_constraint_max, monotone_constraint) : + CUDALeafSplits::GetSplitGains( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output); // gain with split is worse than without split if (current_gain > min_gain_shift) { local_gain = current_gain - min_gain_shift; @@ -1225,10 +1295,16 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( if (sum_left_hessian >= min_sum_hessian_in_leaf && left_count >= min_data_in_leaf && sum_right_hessian >= min_sum_hessian_in_leaf && right_count >= min_data_in_leaf && (!USE_RAND || static_cast(bin + task->mfb_offset) == rand_threshold)) { - double current_gain = CUDALeafSplits::GetSplitGains( - sum_left_gradient, sum_left_hessian, sum_right_gradient, - sum_right_hessian, lambda_l1, - lambda_l2, path_smooth, left_count, right_count, parent_output); + double current_gain = USE_MC ? + CUDALeafSplits::GetSplitGainsMC( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output, + leaf_constraint_min, leaf_constraint_max, monotone_constraint) : + CUDALeafSplits::GetSplitGains( + sum_left_gradient, sum_left_hessian, sum_right_gradient, + sum_right_hessian, lambda_l1, + lambda_l2, path_smooth, left_count, right_count, parent_output); // gain with split is worse than without split if (current_gain > min_gain_shift) { local_gain = current_gain - min_gain_shift; @@ -1259,10 +1335,25 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( const double sum_left_gradient = sum_gradients - sum_right_gradient; const double sum_left_hessian = sum_hessians - sum_right_hessian - kEpsilon; const data_size_t left_count = num_data - right_count; - const double left_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); - const double right_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + // Unconstrained outputs first. The leaf VALUE is clamped into the constraint + // range (MC), but the leaf GAIN stored for future splits must stay + // unconstrained: it becomes the child's parent_gain / min_gain_shift baseline, + // and the CPU reference (BeforeNumerical) recomputes that baseline as the + // unconstrained GetLeafGain of the leaf's sums. + const double left_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, + sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); + const double right_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, + sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + const double left_output = USE_MC ? + (left_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (left_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : left_output_unconstrained)) : + left_output_unconstrained; + const double right_output = USE_MC ? + (right_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (right_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : right_output_unconstrained)) : + right_output_unconstrained; cuda_best_split_info->left_sum_gradients = sum_left_gradient; cuda_best_split_info->left_sum_hessians = sum_left_hessian; cuda_best_split_info->left_count = left_count; @@ -1271,10 +1362,10 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( cuda_best_split_info->right_count = right_count; cuda_best_split_info->left_value = left_output; cuda_best_split_info->left_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, left_output); + sum_left_hessian, lambda_l1, lambda_l2, left_output_unconstrained); cuda_best_split_info->right_value = right_output; cuda_best_split_info->right_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, right_output); + sum_right_hessian, lambda_l1, lambda_l2, right_output_unconstrained); } else { const unsigned int best_bin = (task->na_as_missing && task->mfb_offset == 1) ? threshold_value : static_cast(threshold_value - task->mfb_offset); @@ -1284,10 +1375,25 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( const double sum_right_gradient = sum_gradients - sum_left_gradient; const double sum_right_hessian = sum_hessians - sum_left_hessian - kEpsilon; const data_size_t right_count = num_data - left_count; - const double left_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); - const double right_output = CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + // Unconstrained outputs first. The leaf VALUE is clamped into the constraint + // range (MC), but the leaf GAIN stored for future splits must stay + // unconstrained: it becomes the child's parent_gain / min_gain_shift baseline, + // and the CPU reference (BeforeNumerical) recomputes that baseline as the + // unconstrained GetLeafGain of the leaf's sums. + const double left_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_left_gradient, + sum_left_hessian, lambda_l1, lambda_l2, path_smooth, left_count, parent_output); + const double right_output_unconstrained = + CUDALeafSplits::CalculateSplittedLeafOutput(sum_right_gradient, + sum_right_hessian, lambda_l1, lambda_l2, path_smooth, right_count, parent_output); + const double left_output = USE_MC ? + (left_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (left_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : left_output_unconstrained)) : + left_output_unconstrained; + const double right_output = USE_MC ? + (right_output_unconstrained < leaf_constraint_min ? leaf_constraint_min : + (right_output_unconstrained > leaf_constraint_max ? leaf_constraint_max : right_output_unconstrained)) : + right_output_unconstrained; cuda_best_split_info->left_sum_gradients = sum_left_gradient; cuda_best_split_info->left_sum_hessians = sum_left_hessian; cuda_best_split_info->left_count = left_count; @@ -1296,10 +1402,10 @@ __device__ void FindBestSplitsForLeafKernelInner_GlobalMemory( cuda_best_split_info->right_count = right_count; cuda_best_split_info->left_value = left_output; cuda_best_split_info->left_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_left_gradient, - sum_left_hessian, lambda_l1, lambda_l2, left_output); + sum_left_hessian, lambda_l1, lambda_l2, left_output_unconstrained); cuda_best_split_info->right_value = right_output; cuda_best_split_info->right_gain = CUDALeafSplits::GetLeafGainGivenOutput(sum_right_gradient, - sum_right_hessian, lambda_l1, lambda_l2, right_output); + sum_right_hessian, lambda_l1, lambda_l2, right_output_unconstrained); } } } @@ -1590,7 +1696,7 @@ __device__ void FindBestSplitsForLeafKernelCategoricalInner_GlobalMemory( } } -template +template __global__ void FindBestSplitsForLeafKernel_GlobalMemory( // input feature information const int8_t* is_feature_used_bytree, @@ -1617,11 +1723,18 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( // global num data in leaf const data_size_t global_num_data_in_smaller_leaf, const data_size_t global_num_data_in_larger_leaf, + // monotone leaf constraints + const double smaller_leaf_constraint_min, + const double smaller_leaf_constraint_max, + const double larger_leaf_constraint_min, + const double larger_leaf_constraint_max, // buffer hist_t* feature_hist_grad_buffer, hist_t* feature_hist_hess_buffer, hist_t* feature_hist_stat_buffer, data_size_t* feature_hist_index_buffer) { + const double leaf_constraint_min = IS_LARGER ? larger_leaf_constraint_min : smaller_leaf_constraint_min; + const double leaf_constraint_max = IS_LARGER ? larger_leaf_constraint_max : smaller_leaf_constraint_max; const unsigned int task_index = blockIdx.x; const SplitFindTask* task = tasks + task_index; const double parent_gain = IS_LARGER ? larger_leaf_splits->gain : smaller_leaf_splits->gain; @@ -1673,7 +1786,7 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( out); } else { if (!task->reverse) { - FindBestSplitsForLeafKernelInner_GlobalMemory( + FindBestSplitsForLeafKernelInner_GlobalMemory( // input feature information hist_ptr, // input task information @@ -1692,13 +1805,16 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( sum_hessians, num_data, parent_output, + // monotone constraint information + leaf_constraint_min, + leaf_constraint_max, // output parameters out, // buffer hist_grad_buffer_ptr, hist_hess_buffer_ptr); } else { - FindBestSplitsForLeafKernelInner_GlobalMemory( + FindBestSplitsForLeafKernelInner_GlobalMemory( // input feature information hist_ptr, // input task information @@ -1717,6 +1833,9 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( sum_hessians, num_data, parent_output, + // monotone constraint information + leaf_constraint_min, + leaf_constraint_max, // output parameters out, // buffer @@ -1737,7 +1856,11 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( const bool is_smaller_leaf_valid, \ const bool is_larger_leaf_valid, \ const data_size_t global_num_data_in_smaller_leaf, \ - const data_size_t global_num_data_in_larger_leaf + const data_size_t global_num_data_in_larger_leaf, \ + const double smaller_leaf_constraint_min, \ + const double smaller_leaf_constraint_max, \ + const double larger_leaf_constraint_min, \ + const double larger_leaf_constraint_max #define LaunchFindBestSplitsForLeafKernel_ARGS \ smaller_leaf_splits, \ @@ -1747,7 +1870,11 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( is_smaller_leaf_valid, \ is_larger_leaf_valid, \ global_num_data_in_smaller_leaf, \ - global_num_data_in_larger_leaf + global_num_data_in_larger_leaf, \ + smaller_leaf_constraint_min, \ + smaller_leaf_constraint_max, \ + larger_leaf_constraint_min, \ + larger_leaf_constraint_max #define FindBestSplitsForLeafKernel_ARGS \ num_tasks_, \ @@ -1769,6 +1896,12 @@ __global__ void FindBestSplitsForLeafKernel_GlobalMemory( global_num_data_in_smaller_leaf, \ global_num_data_in_larger_leaf +#define FindBestSplitsForLeafKernel_CONSTRAINT_ARGS \ + smaller_leaf_constraint_min, \ + smaller_leaf_constraint_max, \ + larger_leaf_constraint_min, \ + larger_leaf_constraint_max + #define GlobalMemory_Buffer_ARGS \ cuda_feature_hist_grad_buffer_.RawData(), \ cuda_feature_hist_hess_buffer_.RawData(), \ @@ -1828,16 +1961,28 @@ void CUDABestSplitFinder::LaunchFindBestSplitsForLeafKernelInner2(LaunchFindBest } if (!use_global_memory_) { if (is_smaller_leaf_valid) { - FindBestSplitsForLeafKernel - <<>> - (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS); + if (use_monotone_constraints_) { + FindBestSplitsForLeafKernel + <<>> + (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS); + } else { + FindBestSplitsForLeafKernel + <<>> + (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS); + } } // No device sync here: the larger-leaf launch below waits on subtract_done_event via // its stream, and the smaller/larger leaves write disjoint cuda_best_split_info_ slots. if (is_larger_leaf_valid) { - FindBestSplitsForLeafKernel - <<>> - (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS); + if (use_monotone_constraints_) { + FindBestSplitsForLeafKernel + <<>> + (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS); + } else { + FindBestSplitsForLeafKernel + <<>> + (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS); + } } } else { // Global-memory path (large-dataset fallback, not covered by the shared-memory @@ -1845,21 +1990,34 @@ void CUDABestSplitFinder::LaunchFindBestSplitsForLeafKernelInner2(LaunchFindBest // keep the device sync because the smaller and larger launches share // cuda_feature_hist_grad/hess_buffer_ and must not run concurrently. if (is_smaller_leaf_valid) { - FindBestSplitsForLeafKernel_GlobalMemory - <<>> - (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS, GlobalMemory_Buffer_ARGS); + if (use_monotone_constraints_) { + FindBestSplitsForLeafKernel_GlobalMemory + <<>> + (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS, GlobalMemory_Buffer_ARGS); + } else { + FindBestSplitsForLeafKernel_GlobalMemory + <<>> + (is_feature_used_by_smaller_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS, GlobalMemory_Buffer_ARGS); + } } SynchronizeCUDADevice(__FILE__, __LINE__); if (is_larger_leaf_valid) { - FindBestSplitsForLeafKernel_GlobalMemory - <<>> - (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS, GlobalMemory_Buffer_ARGS); + if (use_monotone_constraints_) { + FindBestSplitsForLeafKernel_GlobalMemory + <<>> + (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS, GlobalMemory_Buffer_ARGS); + } else { + FindBestSplitsForLeafKernel_GlobalMemory + <<>> + (is_feature_used_by_larger_node, FindBestSplitsForLeafKernel_ARGS, FindBestSplitsForLeafKernel_CONSTRAINT_ARGS, GlobalMemory_Buffer_ARGS); + } } } } #undef LaunchFindBestSplitsForLeafKernel_PARAMS #undef FindBestSplitsForLeafKernel_ARGS +#undef FindBestSplitsForLeafKernel_CONSTRAINT_ARGS #undef GlobalMemory_Buffer_ARGS @@ -2034,18 +2192,24 @@ __global__ void FindBestSplitsForLevelKernel( // no categorical branch: the batched path is gated on !has_categorical_feature_ const hist_t* hist_ptr = leaf_splits->hist_in_leaf + task->hist_offset * 2; if (!task->reverse) { - FindBestSplitsForLeafKernelInner( + FindBestSplitsForLeafKernelInner( hist_ptr, task, cuda_random, lambda_l1, lambda_l2, path_smooth, min_data_in_leaf, min_sum_hessian_in_leaf, min_gain_to_split, parent_gain, sum_gradients, sum_hessians, num_data, parent_output, + // hybrid growth is gated off whenever monotone_constraints is set, so the + // batched path never sees a binding constraint: pass identity bounds. + -DBL_MAX, DBL_MAX, out); } else { - FindBestSplitsForLeafKernelInner( + FindBestSplitsForLeafKernelInner( hist_ptr, task, cuda_random, lambda_l1, lambda_l2, path_smooth, min_data_in_leaf, min_sum_hessian_in_leaf, min_gain_to_split, parent_gain, sum_gradients, sum_hessians, num_data, parent_output, + // hybrid growth is gated off whenever monotone_constraints is set, so the + // batched path never sees a binding constraint: pass identity bounds. + -DBL_MAX, DBL_MAX, out); } } else { diff --git a/src/treelearner/cuda/cuda_best_split_finder.hpp b/src/treelearner/cuda/cuda_best_split_finder.hpp index ebf469a21a90..d16fcdd88927 100644 --- a/src/treelearner/cuda/cuda_best_split_finder.hpp +++ b/src/treelearner/cuda/cuda_best_split_finder.hpp @@ -30,6 +30,9 @@ namespace LightGBM { struct SplitFindTask { int inner_feature_index; + // monotone constraint for the (real) feature underlying this task: + // -1 decreasing, +1 increasing, 0 none. Always 0 for categorical tasks. + int8_t monotone_type; bool reverse; bool skip_default_bin; bool na_as_missing; @@ -90,6 +93,10 @@ class CUDABestSplitFinder { const uint8_t larger_num_bits_in_histogram_bins, const bool smaller_leaf_below_max_depth, const bool larger_leaf_below_max_depth, + const double smaller_leaf_constraint_min, + const double smaller_leaf_constraint_max, + const double larger_leaf_constraint_min, + const double larger_leaf_constraint_max, const bool synchronize = true); /*! \brief whether the batched per-level find+sync path supports the current @@ -217,7 +224,11 @@ class CUDABestSplitFinder { const bool is_smaller_leaf_valid, \ const bool is_larger_leaf_valid, \ const data_size_t global_num_data_in_smaller_leaf, \ - const data_size_t global_num_data_in_larger_leaf + const data_size_t global_num_data_in_larger_leaf, \ + const double smaller_leaf_constraint_min, \ + const double smaller_leaf_constraint_max, \ + const double larger_leaf_constraint_min, \ + const double larger_leaf_constraint_max void LaunchFindBestSplitsForLeafKernel(LaunchFindBestSplitsForLeafKernel_PARAMS); @@ -361,6 +372,11 @@ class CUDABestSplitFinder { std::vector is_categorical_; // whether need to select features by node bool select_features_by_node_; + // whether monotone constraints are active (basic method, host-tracked) + bool use_monotone_constraints_; + // per-inner-feature monotone constraint (indexed by inner feature index), + // value taken from config->monotone_constraints[RealFeatureIndex(inner)] + std::vector monotone_constraints_; // CUDA memory, held by this object // for per leaf best split information diff --git a/src/treelearner/cuda/cuda_leaf_splits.hpp b/src/treelearner/cuda/cuda_leaf_splits.hpp index 2f8b368bed0b..b4677d03c2e9 100644 --- a/src/treelearner/cuda/cuda_leaf_splits.hpp +++ b/src/treelearner/cuda/cuda_leaf_splits.hpp @@ -158,6 +158,60 @@ class CUDALeafSplits: public NCCLInfo { l1, l2, path_smooth, right_count, parent_output); } + // Monotone-constraint-aware leaf output: analytic output clamped into the + // per-leaf [constraint_min, constraint_max] interval. Mirrors the CPU + // FeatureHistogram::CalculateSplittedLeafOutput clamp. + template + __device__ static double CalculateSplittedLeafOutputMC(double sum_gradients, + double sum_hessians, double l1, double l2, + double path_smooth, data_size_t num_data, + double parent_output, + double constraint_min, double constraint_max) { + double ret = CalculateSplittedLeafOutput( + sum_gradients, sum_hessians, l1, l2, path_smooth, num_data, parent_output); + if (ret < constraint_min) { + ret = constraint_min; + } else if (ret > constraint_max) { + ret = constraint_max; + } + return ret; + } + + // Monotone-constraint-aware split gain. Mirrors the CPU + // FeatureHistogram::GetSplitGains branch: + // * clamp both child outputs into the per-leaf [min,max] + // * return 0 if the monotonicity relation between children is violated + // * otherwise return GetLeafGainGivenOutput(left)+GetLeafGainGivenOutput(right) + // For a feature with monotone_constraint==0 and an unconstrained leaf + // ([-DBL_MAX, +DBL_MAX]) this is mathematically identical to the non-MC + // analytic gain, and is exactly what the CPU computes when + // monotone_constraints is non-empty (USE_MC is then on for every feature). + template + __device__ static double GetSplitGainsMC(double sum_left_gradients, + double sum_left_hessians, + double sum_right_gradients, + double sum_right_hessians, + double l1, double l2, + double path_smooth, + data_size_t left_count, + data_size_t right_count, + double parent_output, + double constraint_min, double constraint_max, + int8_t monotone_constraint) { + const double left_output = CalculateSplittedLeafOutputMC( + sum_left_gradients, sum_left_hessians, l1, l2, path_smooth, left_count, + parent_output, constraint_min, constraint_max); + const double right_output = CalculateSplittedLeafOutputMC( + sum_right_gradients, sum_right_hessians, l1, l2, path_smooth, right_count, + parent_output, constraint_min, constraint_max); + if (((monotone_constraint > 0) && (left_output > right_output)) || + ((monotone_constraint < 0) && (left_output < right_output))) { + return 0; + } + return GetLeafGainGivenOutput(sum_left_gradients, sum_left_hessians, l1, l2, left_output) + + GetLeafGainGivenOutput(sum_right_gradients, sum_right_hessians, l1, l2, right_output); + } + private: void LaunchInitValuesEmptyKernel(); diff --git a/src/treelearner/cuda/cuda_single_gpu_tree_learner.cpp b/src/treelearner/cuda/cuda_single_gpu_tree_learner.cpp index 79e11314d46c..717c107390b2 100644 --- a/src/treelearner/cuda/cuda_single_gpu_tree_learner.cpp +++ b/src/treelearner/cuda/cuda_single_gpu_tree_learner.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -74,6 +75,15 @@ void CUDASingleGPUTreeLearner::Init(const Dataset* train_data, bool is_constant_ // it (falls back to the classic one-split-at-a-time leaf-wise loop everywhere) const char* hybrid_env = std::getenv("EXABOOST_HYBRID_GROWTH"); use_hybrid_growth_ = (hybrid_env == nullptr || std::string(hybrid_env) != std::string("0")); + // Monotone constraints are inherited down the tree: a leaf's [min, max] bounds + // come from the splits already applied above it. Level-batched growth scores a + // whole level before applying any of it, so the children of a level's splits + // would be searched against their parents' stale bounds. Fall back to the + // classic leaf-wise loop, the same way the batched path is gated off for + // categorical features. + if (!config_->monotone_constraints.empty()) { + use_hybrid_growth_ = false; + } // batched per-level kernels for the hybrid prefix (one construct/fix/subtract/ // find/sync launch per level instead of per pair); "0" keeps the per-pair path const char* batch_env = std::getenv("EXABOOST_HYBRID_BATCH_KERNELS"); @@ -111,6 +121,19 @@ void CUDASingleGPUTreeLearner::Init(const Dataset* train_data, bool is_constant_ leaf_sum_gradients_.resize(config_->num_leaves, 0.0f); leaf_sum_hessians_.resize(config_->num_leaves, 0.0f); + use_monotone_constraints_ = !config_->monotone_constraints.empty(); + if (use_monotone_constraints_) { + leaf_constraint_min_.resize(config_->num_leaves, -std::numeric_limits::max()); + leaf_constraint_max_.resize(config_->num_leaves, std::numeric_limits::max()); + // config_->monotone_constraints is real-indexed; map to inner feature index. + monotone_constraints_.resize(train_data_->num_features(), 0); + for (int inner_feature_index = 0; inner_feature_index < train_data_->num_features(); ++inner_feature_index) { + const int real_feature_index = train_data_->RealFeatureIndex(inner_feature_index); + monotone_constraints_[inner_feature_index] = + config_->monotone_constraints[real_feature_index]; + } + } + if (!boosting_on_cuda_) { cuda_gradients_.Resize(static_cast(num_data_)); cuda_hessians_.Resize(static_cast(num_data_)); @@ -214,6 +237,13 @@ void CUDASingleGPUTreeLearner::BeforeTrain() { smaller_leaf_index_ = 0; larger_leaf_index_ = -1; + if (use_monotone_constraints_) { + std::fill(leaf_constraint_min_.begin(), leaf_constraint_min_.end(), + -std::numeric_limits::max()); + std::fill(leaf_constraint_max_.begin(), leaf_constraint_max_.end(), + std::numeric_limits::max()); + } + if (nccl_communicator_ != nullptr) { leaf_to_hist_index_map_.resize(config_->num_leaves, -1); leaf_to_hist_index_map_[0] = 0; @@ -608,6 +638,15 @@ void CUDASingleGPUTreeLearner::EnqueuePairBestSplitSearch(const CUDATree* tree, larger_num_bits_bin); SelectFeatureByNode(tree); + // monotone constraints for this pair (identity bounds when unused) + const double smaller_leaf_constraint_min = use_monotone_constraints_ ? + leaf_constraint_min_[smaller_leaf_index] : -std::numeric_limits::max(); + const double smaller_leaf_constraint_max = use_monotone_constraints_ ? + leaf_constraint_max_[smaller_leaf_index] : std::numeric_limits::max(); + const double larger_leaf_constraint_min = (use_monotone_constraints_ && larger_leaf_index >= 0) ? + leaf_constraint_min_[larger_leaf_index] : -std::numeric_limits::max(); + const double larger_leaf_constraint_max = (use_monotone_constraints_ && larger_leaf_index >= 0) ? + leaf_constraint_max_[larger_leaf_index] : std::numeric_limits::max(); if (config_->use_quantized_grad) { const uint8_t smaller_leaf_num_bits_bin = nccl_communicator_ == nullptr ? @@ -629,6 +668,8 @@ void CUDASingleGPUTreeLearner::EnqueuePairBestSplitSearch(const CUDATree* tree, config_->max_depth <= 0 || GrowthLeafDepth(tree, smaller_leaf_index) < config_->max_depth, larger_leaf_index < 0 || config_->max_depth <= 0 || GrowthLeafDepth(tree, larger_leaf_index) < config_->max_depth, + smaller_leaf_constraint_min, smaller_leaf_constraint_max, + larger_leaf_constraint_min, larger_leaf_constraint_max, synchronize); } else { cuda_best_split_finder_->FindBestSplitsForLeaf( @@ -641,6 +682,8 @@ void CUDASingleGPUTreeLearner::EnqueuePairBestSplitSearch(const CUDATree* tree, config_->max_depth <= 0 || GrowthLeafDepth(tree, smaller_leaf_index) < config_->max_depth, larger_leaf_index < 0 || config_->max_depth <= 0 || GrowthLeafDepth(tree, larger_leaf_index) < config_->max_depth, + smaller_leaf_constraint_min, smaller_leaf_constraint_max, + larger_leaf_constraint_min, larger_leaf_constraint_max, synchronize); } global_timer.Stop("CUDASingleGPUTreeLearner::FindBestSplitsForLeaf"); @@ -1758,6 +1801,18 @@ Tree* CUDASingleGPUTreeLearner::Train(const score_t* gradients, const int right_leaf_index = ApplySplit(tree.get(), best_split_info, best_leaf_index_); + if (use_monotone_constraints_) { + const int split_inner_feature = leaf_best_split_feature_[best_leaf_index_]; + const bool is_numerical_split = + train_data_->FeatureBinMapper(split_inner_feature)->bin_type() == BinType::NumericalBin; + double host_left_value = 0.0; + double host_right_value = 0.0; + CopyFromCUDADeviceToHost(&host_left_value, &best_split_info->left_value, 1, __FILE__, __LINE__); + CopyFromCUDADeviceToHost(&host_right_value, &best_split_info->right_value, 1, __FILE__, __LINE__); + UpdateLeafConstraints(best_leaf_index_, right_leaf_index, split_inner_feature, + host_left_value, host_right_value, is_numerical_split); + } + if (nccl_communicator_ != nullptr) { smaller_leaf_index_ = (global_num_data_in_leaf_[best_leaf_index_] < global_num_data_in_leaf_[right_leaf_index] ? best_leaf_index_ : right_leaf_index); larger_leaf_index_ = (smaller_leaf_index_ == best_leaf_index_ ? right_leaf_index : best_leaf_index_); @@ -1920,6 +1975,16 @@ int CUDASingleGPUTreeLearner::ForceSplitsCUDA(CUDATree* tree, int* num_splits_do false, 0, 0, 0); // regular best-split search for the active pair: populates the device per-leaf cache SelectFeatureByNode(tree); + // Same monotone constraints the main search applies; without them the splits + // cached here (and later picked by FindBestFromAllSplits) could violate them. + const double forced_smaller_constraint_min = use_monotone_constraints_ ? + leaf_constraint_min_[smaller_leaf_index_] : -std::numeric_limits::max(); + const double forced_smaller_constraint_max = use_monotone_constraints_ ? + leaf_constraint_max_[smaller_leaf_index_] : std::numeric_limits::max(); + const double forced_larger_constraint_min = (use_monotone_constraints_ && larger_leaf_index_ >= 0) ? + leaf_constraint_min_[larger_leaf_index_] : -std::numeric_limits::max(); + const double forced_larger_constraint_max = (use_monotone_constraints_ && larger_leaf_index_ >= 0) ? + leaf_constraint_max_[larger_leaf_index_] : std::numeric_limits::max(); cuda_best_split_finder_->FindBestSplitsForLeaf( cuda_smaller_leaf_splits_->GetCUDAStruct(), cuda_larger_leaf_splits_->GetCUDAStruct(), @@ -1929,7 +1994,9 @@ int CUDASingleGPUTreeLearner::ForceSplitsCUDA(CUDATree* tree, int* num_splits_do nullptr, nullptr, 0, 0, config_->max_depth <= 0 || tree->leaf_depth(smaller_leaf_index_) < config_->max_depth, larger_leaf_index_ < 0 || config_->max_depth <= 0 || - tree->leaf_depth(larger_leaf_index_) < config_->max_depth); + tree->leaf_depth(larger_leaf_index_) < config_->max_depth, + forced_smaller_constraint_min, forced_smaller_constraint_max, + forced_larger_constraint_min, forced_larger_constraint_max); // sync host-side best-split arrays with the device cache for the searched pair, so // the main loop's FindBestFromAllSplits can pick these leaves later with consistent // host (feature, threshold) and device (gain, sums) information @@ -2031,6 +2098,34 @@ int CUDASingleGPUTreeLearner::ForceSplitsCUDA(CUDATree* tree, int* num_splits_do return result_count; } +void CUDASingleGPUTreeLearner::UpdateLeafConstraints( + const int left_leaf, const int right_leaf, const int inner_feature_index, + const double left_value, const double right_value, const bool is_numerical_split) { + // Mirror of BasicLeafConstraints::Update (monotone_constraints.hpp). The new leaf + // inherits the parent's [min,max], then, for a numerical split on a monotone + // feature, the mid-point of the (clamped) child outputs becomes a min/max bound + // for the two children. + leaf_constraint_min_[right_leaf] = leaf_constraint_min_[left_leaf]; + leaf_constraint_max_[right_leaf] = leaf_constraint_max_[left_leaf]; + if (!is_numerical_split) { + return; + } + const int8_t monotone_type = monotone_constraints_[inner_feature_index]; + if (monotone_type == 0) { + return; + } + const double mid = (left_value + right_value) / 2.0; + if (monotone_type < 0) { + // decreasing: left child (parent leaf) gets a min bound, right child a max bound + leaf_constraint_min_[left_leaf] = std::max(mid, leaf_constraint_min_[left_leaf]); + leaf_constraint_max_[right_leaf] = std::min(mid, leaf_constraint_max_[right_leaf]); + } else { + // increasing: left child gets a max bound, right child a min bound + leaf_constraint_max_[left_leaf] = std::min(mid, leaf_constraint_max_[left_leaf]); + leaf_constraint_min_[right_leaf] = std::max(mid, leaf_constraint_min_[right_leaf]); + } +} + void CUDASingleGPUTreeLearner::ResetTrainingData( const Dataset* train_data, bool is_constant_hessian) { diff --git a/src/treelearner/cuda/cuda_single_gpu_tree_learner.hpp b/src/treelearner/cuda/cuda_single_gpu_tree_learner.hpp index 459af004ff47..d9c2a0fbaee9 100644 --- a/src/treelearner/cuda/cuda_single_gpu_tree_learner.hpp +++ b/src/treelearner/cuda/cuda_single_gpu_tree_learner.hpp @@ -339,6 +339,14 @@ class CUDASingleGPUTreeLearner: public SerialTreeLearner, public NCCLInfo { const uint8_t* compact_gather_src_ = nullptr; bool compact_gather_src_is_4bit_ = false; + // Mirror of BasicLeafConstraints::Update for the CUDA path. Given a just-applied + // numerical split on `inner_feature_index`, the parent leaf (`left_leaf`) and the + // newly created leaf (`right_leaf`), and the (clamped) child outputs, update the + // host-side per-leaf [min,max] constraint arrays. + void UpdateLeafConstraints( + const int left_leaf, const int right_leaf, const int inner_feature_index, + const double left_value, const double right_value, const bool is_numerical_split); + // number of threads on CPU int num_threads_; @@ -417,6 +425,14 @@ class CUDASingleGPUTreeLearner: public SerialTreeLearner, public NCCLInfo { std::vector leaf_data_start_; std::vector leaf_sum_gradients_; std::vector leaf_sum_hessians_; + // host-side per-leaf monotone constraints (basic method). [-DBL_MAX, +DBL_MAX] + // means unconstrained. Empty/unused when monotone_constraints is not set. + std::vector leaf_constraint_min_; + std::vector leaf_constraint_max_; + // per-inner-feature monotone constraint, mapped from the real-indexed + // config_->monotone_constraints; empty when monotone_constraints is not set. + std::vector monotone_constraints_; + bool use_monotone_constraints_; int smaller_leaf_index_; int larger_leaf_index_; int best_leaf_index_; diff --git a/tests/python_package_test/test_dual.py b/tests/python_package_test/test_dual.py index c910a4318260..da29f5d57ae5 100644 --- a/tests/python_package_test/test_dual.py +++ b/tests/python_package_test/test_dual.py @@ -968,3 +968,129 @@ def test_cuda_linear_tree_handles_nan_like_cpu(): preds[device_type] = lgb.train(params, ds, num_boost_round=30).predict(X) max_diff = float(np.max(np.abs(preds["cpu"] - preds["cuda"]))) assert max_diff < 1e-6, f"CUDA linear tree (NaN) diverges from CPU: max|diff|={max_diff:.3e}" + + +def _train_monotone(device_type, constraints, num_boost_round, seed=0): + rng = np.random.RandomState(seed) + X = rng.rand(600, 3) + y = 5 * X[:, 0] - 5 * X[:, 1] + 0.7 * X[:, 2] + 0.5 * rng.rand(600) + params = { + "objective": "regression", + "monotone_constraints": constraints, + "monotone_constraints_method": "basic", + "num_leaves": 31, + "min_data_in_leaf": 5, + "learning_rate": 0.1, + "verbose": -1, + "deterministic": True, + "num_threads": 1, + "seed": 0, + "gpu_use_dp": True, + "force_col_wise": True, + "feature_pre_filter": False, + "device_type": device_type, + } + ds = lgb.Dataset(X, label=y, params={"verbose": -1, "feature_pre_filter": False}) + return lgb.train(params, ds, num_boost_round=num_boost_round), X, y + + +def _monotonicity_violations(bst, constraints, n_grid=500): + """Count monotonicity violations of bst on grid sweeps of each constrained feature.""" + count = 0 + worst = 0.0 + grid = np.linspace(0, 1, n_grid) + for j, c in enumerate(constraints): + if c == 0: + continue + for base_value in (0.2, 0.5, 0.8): + base = np.full(3, base_value) + G = np.tile(base, (n_grid, 1)) + G[:, j] = grid + diffs = np.diff(bst.predict(G)) + bad = -diffs if c > 0 else diffs + violations = bad[bad > 1e-12] + count += len(violations) + if len(violations): + worst = max(worst, float(violations.max())) + return count, worst + + +@_REQUIRES_CUDA +@pytest.mark.parametrize("constraints", [[1, -1, 0], [1, 1, 1], [-1, 0, 1], [1, 0, 0]]) +@pytest.mark.parametrize("num_boost_round", [1, 30, 100]) +def test_cuda_monotone_constraints_are_enforced(constraints, num_boost_round): + """CUDA must enforce basic-method monotone constraints: zero violations. + + Regression test for monotone_constraints being silently ignored on CUDA: the + best-split kernels had no constraint plumbing, so CUDA produced models that + violated the requested monotonic relationships (up to 57 violations of + magnitude 0.68 on this data before the fix). + """ + bst, _, _ = _train_monotone("cuda", constraints, num_boost_round) + count, worst = _monotonicity_violations(bst, constraints) + assert count == 0, f"CUDA model violates monotone constraints {constraints}: {count} violations, worst={worst:.3e}" + + +@_REQUIRES_CUDA +@pytest.mark.parametrize("constraints", [[1, -1, 0], [1, 1, 1], [-1, 0, 1]]) +def test_cuda_monotone_constraints_match_cpu_quality(constraints): + """CUDA monotone training quality must match CPU. + + CPU and CUDA may build different (both valid) constrained trees because CPU + additionally prunes features via its is_splittable cache, so exact tree + equality is not required. But both must (a) enforce the constraints and + (b) reach equivalent training loss (within 5%). + """ + num_boost_round = 50 + bst_cpu, X, y = _train_monotone("cpu", constraints, num_boost_round) + bst_cuda, _, _ = _train_monotone("cuda", constraints, num_boost_round) + + # both enforce + for bst, name in ((bst_cpu, "cpu"), (bst_cuda, "cuda")): + count, worst = _monotonicity_violations(bst, constraints) + assert count == 0, f"{name} violates constraints: {count} violations, worst={worst:.3e}" + + # equivalent quality + mse_cpu = float(np.mean((y - bst_cpu.predict(X)) ** 2)) + mse_cuda = float(np.mean((y - bst_cuda.predict(X)) ** 2)) + assert mse_cuda <= mse_cpu * 1.05, f"CUDA mse {mse_cuda} much worse than CPU mse {mse_cpu}" + + +@_REQUIRES_CUDA +@pytest.mark.parametrize("num_boost_round", [1, 30]) +def test_cuda_monotone_noop_constraints_match_cpu_exactly(num_boost_round): + """With all-zero constraints the MC code path must be a no-op: + predictions must match CPU bit-for-bit.""" + bst_cpu, X, _ = _train_monotone("cpu", [0, 0, 0], num_boost_round) + bst_cuda, _, _ = _train_monotone("cuda", [0, 0, 0], num_boost_round) + np.testing.assert_allclose( + bst_cpu.predict(X), + bst_cuda.predict(X), + rtol=0, + atol=1e-10, + err_msg="all-zero monotone constraints must not change CUDA results", + ) + + +@_REQUIRES_CUDA +def test_cuda_monotone_unsupported_configs_raise(): + """Only the basic method, monotone_penalty=0, and full-precision training are + supported on CUDA; other monotone configurations must be rejected loudly.""" + rng = np.random.RandomState(0) + X = rng.rand(200, 3) + y = X[:, 0] - X[:, 1] + base = { + "objective": "regression", + "monotone_constraints": [1, -1, 0], + "device_type": "cuda", + "num_leaves": 7, + "verbose": -1, + } + for bad, expected in ( + ({"monotone_constraints_method": "intermediate"}, r'only supports the "basic" monotone_constraints_method'), + ({"monotone_constraints_method": "advanced"}, r'only supports the "basic" monotone_constraints_method'), + ({"monotone_penalty": 1.0}, r"monotone_penalty is not supported with device_type=cuda"), + ({"use_quantized_grad": True}, r"monotone_constraints is not supported with use_quantized_grad"), + ): + with pytest.raises(lgb.basic.LightGBMError, match=expected): + lgb.train({**base, **bad}, lgb.Dataset(X, label=y, params={"verbose": -1}), num_boost_round=1)