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)