diff --git a/cpp/include/cuml/tree/decisiontree.hpp b/cpp/include/cuml/tree/decisiontree.hpp index abbcf7b1e6..747f091ddb 100644 --- a/cpp/include/cuml/tree/decisiontree.hpp +++ b/cpp/include/cuml/tree/decisiontree.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,6 +10,7 @@ #include +#include #include #include diff --git a/cpp/include/cuml/tree/flatnode.h b/cpp/include/cuml/tree/flatnode.h index 6cfaf840ac..72efab6c01 100644 --- a/cpp/include/cuml/tree/flatnode.h +++ b/cpp/include/cuml/tree/flatnode.h @@ -1,10 +1,12 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2021, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once +#include + // We want to define some functions as usable on device // But need to guard against this file being compiled by a host compiler #ifdef __CUDACC__ @@ -17,18 +19,20 @@ * A node in Decision Tree. * @tparam T data type * @tparam L label type - * @tparam IdxT type used for indexing operations */ -template +template struct SparseTreeNode { private: - IdxT colid = 0; - DataT quesval = DataT(0); - DataT best_metric_val = DataT(0); - IdxT left_child_id = -1; - IdxT instance_count = 0; - FLATNODE_HD SparseTreeNode( - IdxT colid, DataT quesval, DataT best_metric_val, int64_t left_child_id, IdxT instance_count) + std::int64_t colid = 0; + DataT quesval = DataT(0); + DataT best_metric_val = DataT(0); + std::int64_t left_child_id = -1; + std::int64_t instance_count = 0; + FLATNODE_HD SparseTreeNode(std::int64_t colid, + DataT quesval, + DataT best_metric_val, + std::int64_t left_child_id, + std::int64_t instance_count) : colid(colid), quesval(quesval), best_metric_val(best_metric_val), @@ -38,22 +42,24 @@ struct SparseTreeNode { } public: - FLATNODE_HD IdxT ColumnId() const { return colid; } + FLATNODE_HD std::int64_t ColumnId() const { return colid; } FLATNODE_HD DataT QueryValue() const { return quesval; } FLATNODE_HD DataT BestMetric() const { return best_metric_val; } - FLATNODE_HD int64_t LeftChildId() const { return left_child_id; } - FLATNODE_HD int64_t RightChildId() const { return left_child_id + 1; } - FLATNODE_HD IdxT InstanceCount() const { return instance_count; } + FLATNODE_HD std::int64_t LeftChildId() const { return left_child_id; } + FLATNODE_HD std::int64_t RightChildId() const { return left_child_id + 1; } + FLATNODE_HD std::int64_t InstanceCount() const { return instance_count; } - FLATNODE_HD static SparseTreeNode CreateSplitNode( - IdxT colid, DataT quesval, DataT best_metric_val, int64_t left_child_id, IdxT instance_count) + FLATNODE_HD static SparseTreeNode CreateSplitNode(std::int64_t colid, + DataT quesval, + DataT best_metric_val, + std::int64_t left_child_id, + std::int64_t instance_count) { - return SparseTreeNode{ - colid, quesval, best_metric_val, left_child_id, instance_count}; + return SparseTreeNode{colid, quesval, best_metric_val, left_child_id, instance_count}; } - FLATNODE_HD static SparseTreeNode CreateLeafNode(IdxT instance_count) + FLATNODE_HD static SparseTreeNode CreateLeafNode(std::int64_t instance_count) { - return SparseTreeNode{0, 0, 0, -1, instance_count}; + return SparseTreeNode{0, 0, 0, -1, instance_count}; } FLATNODE_HD bool IsLeaf() const { return left_child_id == -1; } bool operator==(const SparseTreeNode& other) const diff --git a/cpp/src/decisiontree/batched-levelalgo/builder.cuh b/cpp/src/decisiontree/batched-levelalgo/builder.cuh index 50573e3d1b..471e7cc80a 100644 --- a/cpp/src/decisiontree/batched-levelalgo/builder.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/builder.cuh @@ -81,7 +81,7 @@ class NodeQueue { bool IsExpandable(const NodeT& n, int depth) { if (depth >= params.max_depth) return false; - if (int(n.InstanceCount()) < params.min_samples_split) return false; + if (n.InstanceCount() < params.min_samples_split) return false; if (params.max_leaves != -1 && tree->leaf_counter >= params.max_leaves) return false; return true; } @@ -147,12 +147,11 @@ template struct Builder { typedef typename ObjectiveT::DataT DataT; typedef typename ObjectiveT::LabelT LabelT; - typedef typename ObjectiveT::IdxT IdxT; typedef typename ObjectiveT::BinT BinT; - typedef SparseTreeNode NodeT; - typedef Split SplitT; - typedef Dataset DatasetT; - typedef Quantiles QuantilesT; + typedef SparseTreeNode NodeT; + typedef Split SplitT; + typedef Dataset DatasetT; + typedef Quantiles QuantilesT; /** default threads per block for most kernels in here */ static constexpr int TPB_DEFAULT = 128; @@ -172,11 +171,11 @@ struct Builder { /** quantiles */ QuantilesT quantiles; /** Tree index */ - IdxT treeid; + std::int64_t treeid; /** Seed used for randomization */ uint64_t seed; /** number of nodes created in the current batch */ - IdxT* n_nodes; + std::int64_t* n_nodes; /** buffer of segmented histograms*/ BinT* histograms; /** mutex array used for atomically updating best split */ @@ -186,9 +185,9 @@ struct Builder { /** current batch of nodes */ NodeWorkItem* d_work_items; /** device AOS to map CTA blocks along dimx to nodes of a batch */ - WorkloadInfo* workload_info; + WorkloadInfo* workload_info; /** host AOS to map CTA blocks along dimx to nodes of a batch */ - WorkloadInfo* h_workload_info; + WorkloadInfo* h_workload_info; /** maximum CTA blocks along dimx */ int max_blocks_dimx = 0; /** host array of splits */ @@ -199,9 +198,9 @@ struct Builder { int n_blks_for_cols = 10; /** Memory alignment value */ const size_t align_value = 512; - IdxT* column_samples; + std::int64_t* column_samples; /** temporary row IDs for row-wise out-of-place partitioning */ - IdxT* partition_row_ids; + std::int64_t* partition_row_ids; /** rmm device workspace buffer */ rmm::device_uvector d_buff; /** pinned host buffer to store the trained nodes */ @@ -211,16 +210,16 @@ struct Builder { Builder(const raft::handle_t& handle, cudaStream_t s, - IdxT treeid, + std::int64_t treeid, uint64_t seed, const DecisionTreeParams& p, const DataT* data, const LabelT* labels, const double* sample_weight, - IdxT n_rows, - IdxT n_cols, - rmm::device_uvector* row_ids, - IdxT n_classes, + std::int64_t n_rows, + std::int64_t n_cols, + rmm::device_uvector* row_ids, + int n_classes, const QuantilesT& q, bool row_major = false) : handle(handle), @@ -233,17 +232,18 @@ struct Builder { sample_weight, n_rows, n_cols, - row_major ? n_cols : IdxT{1}, - row_major ? IdxT{1} : n_rows, - int(row_ids->size()), - max(1, IdxT(params.max_features * n_cols)), + row_major ? n_cols : std::int64_t{1}, + row_major ? std::int64_t{1} : n_rows, + ML::narrow_cast(row_ids->size()), + std::max(std::int64_t{1}, std::int64_t(params.max_features * n_cols)), row_ids->data(), n_classes}, quantiles(q), d_buff(0, builder_stream), distributed(raft::resource::comms_initialized(handle) && handle.get_comms().get_size() > 1) { - max_blocks_dimx = 1 + params.max_batch_size + dataset.n_sampled_rows / TPB_DEFAULT; + max_blocks_dimx = ML::narrow_cast(ML::checked_add( + 1, params.max_batch_size, dataset.n_sampled_rows / TPB_DEFAULT)); ASSERT(q.quantiles_array != nullptr && q.n_bins_array != nullptr, "Currently quantiles need to be computed before this call!"); ASSERT(n_classes >= 1, "n_classes should be at least 1"); @@ -299,26 +299,32 @@ struct Builder { { size_t d_wsize = 0, h_wsize = 0; raft::common::nvtx::range fun_scope("Builder::workspaceSize @builder.cuh [batched-levelalgo]"); - auto max_batch = params.max_batch_size; - size_t max_len_histograms = - max_batch * params.max_n_bins * n_blks_for_cols * dataset.num_outputs; - - d_wsize += calculateAlignedBytes(sizeof(IdxT)); // n_nodes - d_wsize += calculateAlignedBytes(sizeof(BinT) * max_len_histograms); // histograms - d_wsize += calculateAlignedBytes(sizeof(int) * max_batch); // mutex - d_wsize += calculateAlignedBytes(sizeof(SplitT) * max_batch); // splits - d_wsize += calculateAlignedBytes(sizeof(NodeWorkItem) * max_batch); // d_work_Items - d_wsize += // workload_info - calculateAlignedBytes(sizeof(WorkloadInfo) * max_blocks_dimx); - d_wsize += - calculateAlignedBytes(sizeof(IdxT) * max_batch * dataset.n_sampled_cols); // column_samples - d_wsize += calculateAlignedBytes(sizeof(IdxT) * dataset.n_sampled_rows); // partition row IDs + auto max_batch = params.max_batch_size; + size_t max_len_histograms = ML::checked_mul( + max_batch, params.max_n_bins, n_blks_for_cols, dataset.num_outputs); + auto histograms_bytes = ML::checked_mul(sizeof(BinT), max_len_histograms); + auto mutex_bytes = ML::checked_mul(sizeof(int), max_batch); + auto splits_bytes = ML::checked_mul(sizeof(SplitT), max_batch); + auto work_items_bytes = ML::checked_mul(sizeof(NodeWorkItem), max_batch); + auto workload_info_bytes = ML::checked_mul(sizeof(WorkloadInfo), max_blocks_dimx); + auto column_samples_bytes = + ML::checked_mul(sizeof(std::int64_t), max_batch, dataset.n_sampled_cols); + auto partition_row_ids_bytes = + ML::checked_mul(sizeof(std::int64_t), dataset.n_sampled_rows); + + d_wsize += calculateAlignedBytes(sizeof(std::int64_t)); // n_nodes + d_wsize += calculateAlignedBytes(histograms_bytes); // histograms + d_wsize += calculateAlignedBytes(mutex_bytes); // mutex + d_wsize += calculateAlignedBytes(splits_bytes); // splits + d_wsize += calculateAlignedBytes(work_items_bytes); // d_work_Items + d_wsize += calculateAlignedBytes(workload_info_bytes); // workload_info + d_wsize += calculateAlignedBytes(column_samples_bytes); // column_samples + d_wsize += calculateAlignedBytes(partition_row_ids_bytes); // partition row IDs d_wsize += packedHistogramWorkspaceSize(max_len_histograms); // all nodes in the tree - h_wsize += // h_workload_info - calculateAlignedBytes(sizeof(WorkloadInfo) * max_blocks_dimx); - h_wsize += calculateAlignedBytes(sizeof(SplitT) * max_batch); // splits + h_wsize += calculateAlignedBytes(workload_info_bytes); // h_workload_info + h_wsize += calculateAlignedBytes(splits_bytes); // splits return std::make_pair(d_wsize, h_wsize); } @@ -334,36 +340,45 @@ struct Builder { { raft::common::nvtx::range fun_scope( "Builder::assignWorkspace @builder.cuh [batched-levelalgo]"); - auto max_batch = params.max_batch_size; - size_t max_len_histograms = - max_batch * (params.max_n_bins) * n_blks_for_cols * dataset.num_outputs; + auto max_batch = params.max_batch_size; + size_t max_len_histograms = ML::checked_mul( + max_batch, params.max_n_bins, n_blks_for_cols, dataset.num_outputs); + auto histograms_bytes = ML::checked_mul(sizeof(BinT), max_len_histograms); + auto mutex_bytes = ML::checked_mul(sizeof(int), max_batch); + auto splits_bytes = ML::checked_mul(sizeof(SplitT), max_batch); + auto work_items_bytes = ML::checked_mul(sizeof(NodeWorkItem), max_batch); + auto workload_info_bytes = ML::checked_mul(sizeof(WorkloadInfo), max_blocks_dimx); + auto column_samples_bytes = + ML::checked_mul(sizeof(std::int64_t), max_batch, dataset.n_sampled_cols); + auto partition_row_ids_bytes = + ML::checked_mul(sizeof(std::int64_t), dataset.n_sampled_rows); // device - n_nodes = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(IdxT)); + n_nodes = reinterpret_cast(d_wspace); + d_wspace += calculateAlignedBytes(sizeof(std::int64_t)); histograms = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(BinT) * max_len_histograms); + d_wspace += calculateAlignedBytes(histograms_bytes); mutex = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(int) * max_batch); + d_wspace += calculateAlignedBytes(mutex_bytes); splits = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(SplitT) * max_batch); + d_wspace += calculateAlignedBytes(splits_bytes); d_work_items = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(NodeWorkItem) * max_batch); - workload_info = reinterpret_cast*>(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(WorkloadInfo) * max_blocks_dimx); - column_samples = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(IdxT) * max_batch * dataset.n_sampled_cols); - partition_row_ids = reinterpret_cast(d_wspace); - d_wspace += calculateAlignedBytes(sizeof(IdxT) * dataset.n_sampled_rows); + d_wspace += calculateAlignedBytes(work_items_bytes); + workload_info = reinterpret_cast(d_wspace); + d_wspace += calculateAlignedBytes(workload_info_bytes); + column_samples = reinterpret_cast(d_wspace); + d_wspace += calculateAlignedBytes(column_samples_bytes); + partition_row_ids = reinterpret_cast(d_wspace); + d_wspace += calculateAlignedBytes(partition_row_ids_bytes); packed_histograms = reinterpret_cast(d_wspace); d_wspace += packedHistogramWorkspaceSize(max_len_histograms); - RAFT_CUDA_TRY(cudaMemsetAsync(mutex, 0, sizeof(int) * max_batch, builder_stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(mutex, 0, mutex_bytes, builder_stream)); // host - h_workload_info = reinterpret_cast*>(h_wspace); - h_wspace += calculateAlignedBytes(sizeof(WorkloadInfo) * max_blocks_dimx); + h_workload_info = reinterpret_cast(h_wspace); + h_wspace += calculateAlignedBytes(workload_info_bytes); h_splits = reinterpret_cast(h_wspace); - h_wspace += calculateAlignedBytes(sizeof(SplitT) * max_batch); + h_wspace += calculateAlignedBytes(splits_bytes); } /** @@ -389,18 +404,21 @@ struct Builder { } private: - auto updateWorkloadInfo(const std::vector& work_items) + std::size_t updateWorkloadInfo(const std::vector& work_items) { - int n_blocks_dimx = 0; // gridDim.x required for histogram construction + std::size_t n_blocks_dimx = 0; // gridDim.x required for histogram construction for (std::size_t i = 0; i < work_items.size(); i++) { - auto item = work_items[i]; - int n_blocks_per_node = - std::max(raft::ceildiv(item.instances.count, size_t(TPB_DEFAULT)), size_t(1)); - - for (int b = 0; b < n_blocks_per_node; b++) { - h_workload_info[n_blocks_dimx + b] = {int(i), b, n_blocks_per_node}; + auto item = work_items[i]; + auto n_blocks_per_node = std::max( + raft::ceildiv(item.instances.count, std::size_t{TPB_DEFAULT}), std::size_t{1}); + + for (std::size_t b = 0; b < n_blocks_per_node; b++) { + auto workload_idx = ML::checked_add(n_blocks_dimx, b); + h_workload_info[workload_idx] = {ML::narrow_cast(i), + ML::narrow_cast(b), + ML::narrow_cast(n_blocks_per_node)}; } - n_blocks_dimx += n_blocks_per_node; + n_blocks_dimx = ML::checked_add(n_blocks_dimx, n_blocks_per_node); } raft::update_device(workload_info, h_workload_info, n_blocks_dimx, builder_stream); return n_blocks_dimx; @@ -410,13 +428,15 @@ struct Builder { { raft::common::nvtx::range fun_scope("Builder::doSplit @builder.cuh [batched-levelalgo]"); // start fresh on the number of *new* nodes created in this batch - RAFT_CUDA_TRY(cudaMemsetAsync(n_nodes, 0, sizeof(IdxT), builder_stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(n_nodes, 0, sizeof(std::int64_t), builder_stream)); - const IdxT original_n_sampled_cols = dataset.n_sampled_cols; + const std::int64_t original_n_sampled_cols = dataset.n_sampled_cols; ASSERT(original_n_sampled_cols > 0 && original_n_sampled_cols <= dataset.n_cols, "n_sampled_cols must be in [1, n_cols]"); - const std::size_t max_sampling_rounds = - std::size_t((dataset.n_cols + original_n_sampled_cols - 1) / original_n_sampled_cols); + const auto sampling_round_numerator = ML::checked_sub( + ML::checked_add(dataset.n_cols, original_n_sampled_cols), 1); + const auto max_sampling_rounds = ML::narrow_cast( + ML::checked_div(sampling_round_numerator, original_n_sampled_cols)); // The final split chosen for each original work item. Nodes that need // additional feature samples are compacted in active_items, so successful // splits must be copied back to their original batch position. @@ -434,9 +454,10 @@ struct Builder { // Match sklearn's behavior of searching beyond max_features when the // sampled features do not yield a valid split. for (std::size_t round = 0; !active_items.empty() && round < max_sampling_rounds; ++round) { - IdxT sample_offset = IdxT(round) * original_n_sampled_cols; - dataset.n_sampled_cols = - std::min(original_n_sampled_cols, static_cast(dataset.n_cols) - sample_offset); + auto sample_offset = ML::checked_mul(ML::narrow_cast(round), + original_n_sampled_cols); + dataset.n_sampled_cols = std::min( + original_n_sampled_cols, ML::checked_sub(dataset.n_cols, sample_offset)); computeBestSplits(active_items, seed, sample_offset); std::vector retry_items; @@ -458,22 +479,20 @@ struct Builder { // Partition samples once, using the valid split found for each node. Nodes // still without a valid split after all features have been visited remain leaves. - RAFT_CUDA_TRY(cudaMemcpyAsync(splits, - final_splits.data(), - sizeof(SplitT) * work_items.size(), - cudaMemcpyHostToDevice, - builder_stream)); + auto split_copy_bytes = ML::checked_mul(sizeof(SplitT), work_items.size()); + RAFT_CUDA_TRY(cudaMemcpyAsync( + splits, final_splits.data(), split_copy_bytes, cudaMemcpyHostToDevice, builder_stream)); raft::update_device(d_work_items, work_items.data(), work_items.size(), builder_stream); const auto n_partition_blocks = this->updateWorkloadInfo(work_items); raft::common::nvtx::push_range("nodeSplitKernel @builder.cuh [batched-levelalgo]"); - launchNodeSplitKernel(dataset, - d_work_items, - splits, - workload_info, - n_partition_blocks, - work_items.size(), - partition_row_ids, - builder_stream); + launchNodeSplitKernel(dataset, + d_work_items, + splits, + workload_info, + n_partition_blocks, + work_items.size(), + partition_row_ids, + builder_stream); RAFT_CUDA_TRY(cudaPeekAtLastError()); raft::common::nvtx::pop_range(); raft::update_host(h_splits, splits, work_items.size(), builder_stream); @@ -483,17 +502,18 @@ struct Builder { void computeBestSplits(const std::vector& work_items, uint64_t sampling_seed, - IdxT sample_offset) + std::int64_t sample_offset) { - initSplit(splits, work_items.size(), builder_stream); - RAFT_CUDA_TRY(cudaMemsetAsync(mutex, 0, sizeof(int) * params.max_batch_size, builder_stream)); + initSplit(splits, work_items.size(), builder_stream); + auto mutex_bytes = ML::checked_mul(sizeof(int), params.max_batch_size); + RAFT_CUDA_TRY(cudaMemsetAsync(mutex, 0, mutex_bytes, builder_stream)); raft::update_device(d_work_items, work_items.data(), work_items.size(), builder_stream); auto n_blocks_dimx = this->updateWorkloadInfo(work_items); auto split_smem_config = computeSharedMemoryConfig(); sampleFeatures(work_items, sampling_seed, sample_offset); - for (IdxT c = 0; c < dataset.n_sampled_cols; c += n_blks_for_cols) { + for (std::int64_t c = 0; c < dataset.n_sampled_cols; c += n_blks_for_cols) { computeSplit(c, n_blocks_dimx, work_items.size(), split_smem_config); RAFT_CUDA_TRY(cudaPeekAtLastError()); } @@ -503,18 +523,18 @@ struct Builder { void sampleFeatures(const std::vector& work_items, uint64_t sampling_seed, - IdxT sample_offset) + std::int64_t sample_offset) { raft::common::nvtx::range fun_scope("feature-sampling"); - sample_features(column_samples, - d_work_items, - work_items.size(), - treeid, - sampling_seed, - sample_offset, - static_cast(dataset.n_cols), - dataset.n_sampled_cols, - builder_stream); + sample_features(column_samples, + d_work_items, + work_items.size(), + treeid, + sampling_seed, + sample_offset, + dataset.n_cols, + dataset.n_sampled_cols, + builder_stream); RAFT_CUDA_TRY(cudaPeekAtLastError()); } @@ -566,7 +586,7 @@ struct Builder { RAFT_CUDA_TRY(cudaPeekAtLastError()); } - void computeSplit(IdxT col, + void computeSplit(std::int64_t col, size_t n_blocks_dimx, size_t n_work_items, const SharedMemoryConfig& split_smem_config) @@ -577,7 +597,11 @@ struct Builder { auto n_bins = params.max_n_bins; auto n_classes = dataset.num_outputs; // if columns left to be processed lesser than `n_blks_for_cols`, shrink the blocks along dimy - auto n_blocks_dimy = std::min(n_blks_for_cols, dataset.n_sampled_cols - col); + auto remaining_sampled_cols = dataset.n_sampled_cols - col; + auto n_blocks_dimy = n_blks_for_cols; + if (remaining_sampled_cols < n_blocks_dimy) { + n_blocks_dimy = ML::narrow_cast(remaining_sampled_cols); + } dim3 histogram_grid(ML::narrow_cast(n_blocks_dimx), ML::narrow_cast(n_blocks_dimy), 1); @@ -594,63 +618,65 @@ struct Builder { params.split_criterion, params.min_impurity_decrease); raft::common::nvtx::range kernel_scope("computeSplitKernels @builder.cuh [batched-levelalgo]"); - launchBuildHistogramsKernel(histograms, - params.max_n_bins, - dataset, - quantiles, - d_work_items, - col, - column_samples, - objective, - workload_info, - histogram_grid, - split_smem_config, - builder_stream); + launchBuildHistogramsKernel(histograms, + params.max_n_bins, + dataset, + quantiles, + d_work_items, + col, + column_samples, + objective, + workload_info, + histogram_grid, + split_smem_config, + builder_stream); RAFT_CUDA_TRY(cudaPeekAtLastError()); // Distributed RF must aggregate per-rank histograms before split scoring. // The split kernel then sees the same global CDF histogram on every rank. if (distributed) { allReduceHistograms(histograms, len_histograms); } - launchFindBestSplitsKernel(histograms, - params.max_n_bins, - dataset, - quantiles, - col, - column_samples, - mutex, - splits, - objective, - split_grid, - builder_stream); + launchFindBestSplitsKernel(histograms, + params.max_n_bins, + dataset, + quantiles, + col, + column_samples, + mutex, + splits, + objective, + split_grid, + builder_stream); } // Set the leaf value predictions in batch void SetLeafPredictions(std::shared_ptr> tree, const std::vector& instance_ranges) { - tree->vector_leaf.resize(tree->sparsetree.size() * dataset.num_outputs); + auto vector_leaf_size = + ML::checked_mul(tree->sparsetree.size(), dataset.num_outputs); + tree->vector_leaf.resize(vector_leaf_size); ASSERT(tree->sparsetree.size() == instance_ranges.size(), "Expected instance range for each node"); // do this in batch to reduce peak memory usage in extreme cases - std::size_t max_batch_size = min(std::size_t(100000), tree->sparsetree.size()); + std::size_t max_batch_size = min(std::size_t{100000}, tree->sparsetree.size()); + auto max_leaf_values = ML::checked_mul(max_batch_size, dataset.num_outputs); rmm::device_uvector d_tree(max_batch_size, builder_stream); rmm::device_uvector d_instance_ranges(max_batch_size, builder_stream); - rmm::device_uvector d_leaves(max_batch_size * dataset.num_outputs, builder_stream); + rmm::device_uvector d_leaves(max_leaf_values, builder_stream); ObjectiveT objective(dataset.num_outputs, params.min_samples_leaf, params.split_criterion); for (std::size_t batch_begin = 0; batch_begin < tree->sparsetree.size(); batch_begin += max_batch_size) { - std::size_t batch_end = min(batch_begin + max_batch_size, tree->sparsetree.size()); - std::size_t batch_size = batch_end - batch_begin; + std::size_t batch_size = min(max_batch_size, tree->sparsetree.size() - batch_begin); raft::update_device( d_tree.data(), tree->sparsetree.data() + batch_begin, batch_size, builder_stream); raft::update_device( d_instance_ranges.data(), instance_ranges.data() + batch_begin, batch_size, builder_stream); - RAFT_CUDA_TRY( - cudaMemsetAsync(d_leaves.data(), 0, sizeof(DataT) * d_leaves.size(), builder_stream)); - size_t smem_size = sizeof(BinT) * dataset.num_outputs; + auto leaves_bytes = ML::checked_mul(sizeof(DataT), d_leaves.size()); + RAFT_CUDA_TRY(cudaMemsetAsync(d_leaves.data(), 0, leaves_bytes, builder_stream)); + size_t smem_size = ML::checked_mul(sizeof(BinT), dataset.num_outputs); launchLeafKernel(objective, dataset, d_tree.data(), @@ -659,10 +685,10 @@ struct Builder { batch_size, smem_size, builder_stream); - raft::update_host(tree->vector_leaf.data() + batch_begin * dataset.num_outputs, - d_leaves.data(), - batch_size * dataset.num_outputs, - builder_stream); + auto leaf_offset = ML::checked_mul(batch_begin, dataset.num_outputs); + auto leaf_count = ML::checked_mul(batch_size, dataset.num_outputs); + raft::update_host( + tree->vector_leaf.data() + leaf_offset, d_leaves.data(), leaf_count, builder_stream); } } }; // end Builder diff --git a/cpp/src/decisiontree/batched-levelalgo/dataset.h b/cpp/src/decisiontree/batched-levelalgo/dataset.h index c2cc369301..18630a4128 100644 --- a/cpp/src/decisiontree/batched-levelalgo/dataset.h +++ b/cpp/src/decisiontree/batched-levelalgo/dataset.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -12,7 +12,7 @@ namespace ML { namespace DT { -template +template struct Dataset { /** input dataset */ const DataT* data; @@ -29,18 +29,17 @@ struct Dataset { /** column stride in input data elements */ std::int64_t col_stride; /** total sampled rows in dataset */ - IdxT n_sampled_rows; + std::int64_t n_sampled_rows; /** total sampled cols in dataset */ - IdxT n_sampled_cols; + std::int64_t n_sampled_cols; /** indices of sampled rows */ - IdxT* row_ids; + std::int64_t* row_ids; /** Number of classes or regression outputs*/ - IdxT num_outputs; + int num_outputs; - HDI DataT value(IdxT row, IdxT col) const + HDI DataT value(std::int64_t row, std::int64_t col) const { - return data[static_cast(row) * row_stride + - static_cast(col) * col_stride]; + return data[row * row_stride + col * col_stride]; } }; diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh b/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh index 4c9745c35c..68635df491 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh @@ -10,8 +10,10 @@ #include "../quantiles.h" #include "../random_utils.cuh" +#include #include +#include #include #include @@ -43,12 +45,11 @@ struct NodeWorkItem { * This struct has information about workload of a single threadblock of * computeSplit kernels of classification and regression */ -template struct WorkloadInfo { - IdxT nodeid; // Node in the batch on which the threadblock needs to work - IdxT offset_blockid; // Offset threadblock id among all the blocks that are - // working on this node - IdxT num_blocks; // Total number of blocks that are working on the node + std::int64_t nodeid; // Node in the batch on which the threadblock needs to work + std::int64_t offset_blockid; // Offset threadblock id among all the blocks that are + // working on this node + std::int64_t num_blocks; // Total number of blocks that are working on the node }; struct SharedMemoryConfig { @@ -63,44 +64,47 @@ DI OutT* alignPointer(InT dataset) return reinterpret_cast(raft::alignTo(reinterpret_cast(dataset), sizeof(OutT))); } -template -void sample_features(IdxT* column_samples, - const NodeWorkItem* work_items, - size_t work_items_size, - IdxT treeid, - uint64_t seed, - IdxT sample_offset, - IdxT n, - IdxT k, - cudaStream_t stream) +inline void sample_features(std::int64_t* column_samples, + const NodeWorkItem* work_items, + size_t work_items_size, + std::int64_t treeid, + uint64_t seed, + std::int64_t sample_offset, + std::int64_t n, + std::int64_t k, + cudaStream_t stream) { - auto n_column_samples = work_items_size * size_t(k); - auto counting = thrust::make_counting_iterator(0); + RAFT_EXPECTS(k >= 0, "k must be non-negative"); + RAFT_EXPECTS(n >= k, "k must not exceed n"); + + auto sampled_cols = ML::narrow_cast(k); + auto n_column_samples = ML::checked_mul(work_items_size, sampled_cols); + auto counting = thrust::make_counting_iterator(0); thrust::for_each(thrust::cuda::par.on(stream), counting, counting + n_column_samples, - [=] __device__(size_t sample_idx) { - auto node_idx = sample_idx / size_t(k); - IdxT column_index = static_cast(sample_idx % size_t(k)); + [=] __device__(std::size_t sample_idx) { + auto node_idx = sample_idx / sampled_cols; + auto column_index = static_cast(sample_idx % sampled_cols); - const uint32_t nodeid = work_items[node_idx].idx; - uint32_t rng_seed = fnv1a32_hash(seed, treeid, nodeid); + auto nodeid = work_items[node_idx].idx; + uint32_t rng_seed = fnv1a32_hash(seed, treeid, nodeid); - cuda::shuffle_iterator shuffled_features( + cuda::shuffle_iterator shuffled_features( n, cuda::std::minstd_rand(rng_seed), sample_offset); column_samples[sample_idx] = shuffled_features[column_index]; }); } -template -void launchNodeSplitKernel(const Dataset& dataset, +template +void launchNodeSplitKernel(const Dataset& dataset, const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, + Split* splits, + const WorkloadInfo* workload_info, size_t n_blocks_dimx, size_t n_work_items, - IdxT* partition_row_ids, + std::int64_t* partition_row_ids, cudaStream_t builder_stream); template @@ -112,49 +116,29 @@ void launchLeafKernel(ObjectiveT objective, int batch_size, size_t smem_size, cudaStream_t builder_stream); -// Returns the lowest index in `array` whose value is greater or equal to `element`. -// Values outside the quantile range are clamped to the edge bins: values below the -// first quantile return 0, and values above the last quantile return len - 1. -template -HDI IdxT lower_bound(DataT const* array, IdxT len, DataT element) -{ - IdxT start = 0; - IdxT end = len - 1; - IdxT mid; - while (start < end) { - mid = (start + end) / 2; - if (array[mid] < element) { - start = mid + 1; - } else { - end = mid; - } - } - return start; -} - -template +template void launchBuildHistogramsKernel(typename ObjectiveT::BinT* histograms, - IdxT n_bins, - const Dataset& dataset, - const Quantiles& quantiles, + std::int64_t n_bins, + const Dataset& dataset, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); -template +template void launchFindBestSplitsKernel(typename ObjectiveT::BinT* histograms, - IdxT n_bins, - const Dataset& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + std::int64_t n_bins, + const Dataset& dataset, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); @@ -168,7 +152,7 @@ inline void packHistograms(const BinT* in, double* out, std::size_t len, cudaStr { // Counts are packed as doubles so each bin can use one homogeneous arithmetic buffer. This is // exact for current RF problem sizes: integer values up to 2^53 are exactly representable by - // double, and IdxT row indexing is far below that limit. + // double, and RF row indexing is far below that limit. auto op = [in] __device__(double* out, std::size_t i) { auto const bin_idx = i / reduction_buffer_size_v; auto const field = i % reduction_buffer_size_v; diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh b/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh index 0e63275d87..9929f8e73c 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh @@ -16,6 +16,8 @@ #include #include +#include +#include #include #include #include @@ -46,18 +48,18 @@ struct NodeSplitPartitionScanOp { } }; -template -static __global__ void resetLocalLeftCountsKernel(Split* splits, std::size_t n_splits) +template +static __global__ void resetLocalLeftCountsKernel(Split* splits, std::size_t n_splits) { const auto idx = std::size_t(blockIdx.x) * blockDim.x + threadIdx.x; if (idx < n_splits) { splits[idx].local_nLeft = 0; } } -template -static __global__ void countLocalLeftKernel(const Dataset dataset, +template +static __global__ void countLocalLeftKernel(const Dataset dataset, const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info) + Split* splits, + const WorkloadInfo* workload_info) { using BlockReduce = cub::BlockReduce; __shared__ typename BlockReduce::TempStorage temp_storage; @@ -86,13 +88,13 @@ static __global__ void countLocalLeftKernel(const Dataset d // inclusive left count and current row side for each logical row slot in its // node segment; this writer uses that state to place the row into the temporary // partition buffer. -template +template struct NodeSplitPartitionWriter { - Dataset dataset; + Dataset dataset; const NodeWorkItem* work_items; - const Split* splits; - const WorkloadInfo* workload_info; - IdxT* partition_row_ids; + const Split* splits; + const WorkloadInfo* workload_info; + std::int64_t* partition_row_ids; __host__ __device__ void operator()(std::ptrdiff_t index, NodeSplitPartitionState state) const { @@ -118,12 +120,12 @@ struct NodeSplitPartitionWriter { // Copy back only ranges for nodes that actually split. Leaf/invalid nodes keep // their existing row-id order because the scan writer skips them too. -template -static __global__ void nodeSplitCopyBackKernel(const Dataset dataset, +template +static __global__ void nodeSplitCopyBackKernel(const Dataset dataset, const NodeWorkItem* work_items, - const Split* splits, - const WorkloadInfo* workload_info, - const IdxT* partition_row_ids) + const Split* splits, + const WorkloadInfo* workload_info, + const std::int64_t* partition_row_ids) { const auto workload_info_cta = workload_info[blockIdx.x]; const auto nid = workload_info_cta.nodeid; @@ -140,25 +142,25 @@ static __global__ void nodeSplitCopyBackKernel(const Dataset -void launchNodeSplitKernel(const Dataset& dataset, +template +void launchNodeSplitKernel(const Dataset& dataset, const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, + Split* splits, + const WorkloadInfo* workload_info, size_t n_blocks_dimx, size_t n_work_items, - IdxT* partition_row_ids, + std::int64_t* partition_row_ids, cudaStream_t builder_stream) { if (n_blocks_dimx == 0) return; constexpr int reset_tpb = 128; const auto reset_grid = raft::ceildiv(n_work_items, std::size_t{reset_tpb}); - resetLocalLeftCountsKernel + resetLocalLeftCountsKernel <<(reset_grid), reset_tpb, 0, builder_stream>>>( splits, n_work_items); RAFT_CUDA_TRY(cudaPeekAtLastError()); - countLocalLeftKernel + countLocalLeftKernel <<(n_blocks_dimx), TPB, 0, builder_stream>>>( dataset, work_items, splits, workload_info); RAFT_CUDA_TRY(cudaPeekAtLastError()); @@ -195,18 +197,18 @@ void launchNodeSplitKernel(const Dataset& dataset, auto node_keys = thrust::make_transform_iterator(slots_begin, node_key); auto partition_states = thrust::make_transform_iterator(slots_begin, partition_state); auto partition_writer = - thrust::make_tabulate_output_iterator(NodeSplitPartitionWriter{ + thrust::make_tabulate_output_iterator(NodeSplitPartitionWriter{ dataset, work_items, splits, workload_info, partition_row_ids}); thrust::inclusive_scan_by_key(exec_policy, node_keys, node_keys + n_slots, partition_states, partition_writer, - thrust::equal_to{}, + thrust::equal_to{}, NodeSplitPartitionScanOp{}); // The original row_ids buffer remains the source during the scan, so copy back after it finishes. - nodeSplitCopyBackKernel<<>>( + nodeSplitCopyBackKernel<<>>( dataset, work_items, splits, workload_info, partition_row_ids); } @@ -260,8 +262,8 @@ void launchLeafKernel(ObjectiveT objective, * @brief For every threadblock, converts a pdf-histogram to a * cdf-histogram inplace using inclusive block-sum-scan. */ -template -DI BinT pdf_to_cdf(BinT* histogram, IdxT n_bins) +template +DI BinT pdf_to_cdf(BinT* histogram, std::int64_t n_bins) { // Blockscan instance preparation typedef cub::BlockScan BlockScan; @@ -270,7 +272,8 @@ DI BinT pdf_to_cdf(BinT* histogram, IdxT n_bins) // variable to accumulate aggregate of sumscans of previous iterations BinT total_aggregate = BinT(); - for (IdxT tix = threadIdx.x; tix < raft::ceildiv(n_bins, TPB) * TPB; tix += blockDim.x) { + for (std::int64_t tix = threadIdx.x; tix < raft::ceildiv(n_bins, std::int64_t{TPB}) * TPB; + tix += blockDim.x) { BinT result; BinT block_aggregate; BinT element = tix < n_bins ? histogram[tix] : BinT(); @@ -282,33 +285,33 @@ DI BinT pdf_to_cdf(BinT* histogram, IdxT n_bins) return total_aggregate; } -template +template static __global__ void buildHistogramsKernel(typename ObjectiveT::BinT* histograms, - IdxT max_n_bins, - const Dataset dataset, - const Quantiles quantiles, + std::int64_t max_n_bins, + const Dataset dataset, + const Quantiles quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, bool use_global_memory_histogram) { using BinT = typename ObjectiveT::BinT; extern __shared__ char smem[]; - WorkloadInfo workload_info_cta = workload_info[blockIdx.x]; - IdxT nid = workload_info_cta.nodeid; - const auto work_item = work_items[nid]; - auto range_start = work_item.instances.begin; - auto range_len = work_item.instances.count; + WorkloadInfo workload_info_cta = workload_info[blockIdx.x]; + std::int64_t nid = workload_info_cta.nodeid; + const auto work_item = work_items[nid]; + auto range_start = work_item.instances.begin; + auto range_len = work_item.instances.count; - IdxT offset_blockid = workload_info_cta.offset_blockid; - IdxT num_blocks = workload_info_cta.num_blocks; + std::int64_t offset_blockid = workload_info_cta.offset_blockid; + std::int64_t num_blocks = workload_info_cta.num_blocks; - IdxT colIndex = colStart + blockIdx.y; - IdxT col = column_samples[nid * dataset.n_sampled_cols + colIndex]; - int n_bins = quantiles.n_bins_array[col]; + std::int64_t colIndex = colStart + blockIdx.y; + std::int64_t col = column_samples[nid * dataset.n_sampled_cols + colIndex]; + int n_bins = quantiles.n_bins_array[col]; auto n_classes = objective.NumClasses(); auto end = range_start + range_len; @@ -317,17 +320,17 @@ static __global__ void buildHistogramsKernel(typename ObjectiveT::BinT* histogra auto* global_histogram = histograms + histograms_offset; auto* histogram = global_histogram; auto* quantiles_for_split = quantiles.quantiles_array + std::size_t(max_n_bins) * col; - IdxT stride = blockDim.x * num_blocks; - IdxT tid = threadIdx.x + offset_blockid * blockDim.x; + std::int64_t stride = blockDim.x * num_blocks; + std::int64_t tid = threadIdx.x + offset_blockid * blockDim.x; if (!use_global_memory_histogram) { histogram = alignPointer(smem); auto* shared_quantiles = alignPointer(histogram + histogram_len); quantiles_for_split = shared_quantiles; - for (IdxT i = threadIdx.x; i < histogram_len; i += blockDim.x) { + for (std::int64_t i = threadIdx.x; i < histogram_len; i += blockDim.x) { histogram[i] = BinT(); } - for (IdxT b = threadIdx.x; b < n_bins; b += blockDim.x) { + for (std::int64_t b = threadIdx.x; b < n_bins; b += blockDim.x) { shared_quantiles[b] = quantiles.quantiles_array[max_n_bins * col + b]; } __syncthreads(); @@ -338,40 +341,47 @@ static __global__ void buildHistogramsKernel(typename ObjectiveT::BinT* histogra auto data = dataset.value(row, col); auto label = dataset.labels[row]; - IdxT start = lower_bound(quantiles_for_split, n_bins, data); - objective.IncrementHistogram(histogram, n_bins, start, label, dataset, row); + // Search bin indices so lower_bound uses 32-bit distance and advance arithmetic. + auto bin_begin = cuda::counting_iterator(0); + auto bin_end = bin_begin + n_bins; + auto bin_it = ::cuda::std::lower_bound( + bin_begin, bin_end, data, [quantiles_for_split](int bin, DataT value) { + return quantiles_for_split[bin] < value; + }); + auto bin = bin_it == bin_end ? n_bins - 1 : *bin_it; + objective.IncrementHistogram(histogram, n_bins, bin, label, dataset, row); } if (!use_global_memory_histogram) { __syncthreads(); - for (IdxT i = threadIdx.x; i < histogram_len; i += blockDim.x) { + for (std::int64_t i = threadIdx.x; i < histogram_len; i += blockDim.x) { BinT::AtomicAdd(global_histogram + i, histogram[i]); } } } -template +template static __global__ void findBestSplitsKernel(typename ObjectiveT::BinT* histograms, - IdxT max_n_bins, - const Dataset dataset, - const Quantiles quantiles, - IdxT colStart, - const IdxT* column_samples, + std::int64_t max_n_bins, + const Dataset dataset, + const Quantiles quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT objective) { using BinT = typename ObjectiveT::BinT; constexpr int n_split_warps = (TPB + raft::WarpSize - 1) / raft::WarpSize; - __shared__ __align__(alignof(Split)) unsigned char - split_scratch_storage[sizeof(Split) * n_split_warps]; - auto* split_scratch = reinterpret_cast*>(split_scratch_storage); + __shared__ __align__(alignof( + Split)) unsigned char split_scratch_storage[sizeof(Split) * n_split_warps]; + auto* split_scratch = reinterpret_cast*>(split_scratch_storage); - IdxT nid = blockIdx.x; + std::int64_t nid = blockIdx.x; - IdxT colIndex = colStart + blockIdx.y; - IdxT col = column_samples[nid * dataset.n_sampled_cols + colIndex]; - int n_bins = quantiles.n_bins_array[col]; + std::int64_t colIndex = colStart + blockIdx.y; + std::int64_t col = column_samples[nid * dataset.n_sampled_cols + colIndex]; + int n_bins = quantiles.n_bins_array[col]; auto n_classes = objective.NumClasses(); auto histograms_offset = (std::size_t(nid) * gridDim.y + blockIdx.y) * max_n_bins * n_classes; @@ -379,14 +389,14 @@ static __global__ void findBestSplitsKernel(typename ObjectiveT::BinT* histogram auto* quantiles_for_split = quantiles.quantiles_array + std::size_t(max_n_bins) * col; std::int64_t global_sample_count = 0; - for (IdxT c = 0; c < n_classes; ++c) { - global_sample_count += static_cast( - pdf_to_cdf(histogram + n_bins * c, n_bins).Count()); + for (std::int64_t c = 0; c < n_classes; ++c) { + global_sample_count += + static_cast(pdf_to_cdf(histogram + n_bins * c, n_bins).Count()); } __syncthreads(); - Split sp = + Split sp = objective.Gain(histogram, quantiles_for_split, col, global_sample_count, n_bins); __syncthreads(); @@ -394,21 +404,21 @@ static __global__ void findBestSplitsKernel(typename ObjectiveT::BinT* histogram sp.evalBestSplit(split_scratch, splits + nid, mutex + nid, quantiles_for_split, n_bins); } -template +template void launchBuildHistogramsKernel(typename ObjectiveT::BinT* histograms, - IdxT max_n_bins, - const Dataset& dataset, - const Quantiles& quantiles, + std::int64_t max_n_bins, + const Dataset& dataset, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream) { - buildHistogramsKernel + buildHistogramsKernel <<>>( histograms, max_n_bins, @@ -422,29 +432,21 @@ void launchBuildHistogramsKernel(typename ObjectiveT::BinT* histograms, split_smem_config.use_global_memory_histogram); } -template +template void launchFindBestSplitsKernel(typename ObjectiveT::BinT* histograms, - IdxT max_n_bins, - const Dataset& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + std::int64_t max_n_bins, + const Dataset& dataset, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream) { - findBestSplitsKernel - <<>>(histograms, - max_n_bins, - dataset, - quantiles, - colStart, - column_samples, - mutex, - splits, - objective); + findBestSplitsKernel<<>>( + histograms, max_n_bins, dataset, quantiles, colStart, column_samples, mutex, splits, objective); } } // namespace DT diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu index 9f4ee85714..233fc6f9ff 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = double; using LabelT = int; -using IdxT = int; -using ObjectiveT = ClassificationObjectiveFunction; +using ObjectiveT = ClassificationObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu index 81dd3a924d..7b75c462fb 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = float; using LabelT = int; -using IdxT = int; -using ObjectiveT = ClassificationObjectiveFunction; +using ObjectiveT = ClassificationObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu index 0120ca652f..2ceed72cdd 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu @@ -9,47 +9,44 @@ namespace ML { namespace DT { // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchNodeSplitKernel( - const Dataset& dataset, - const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, - size_t n_blocks_dimx, - size_t n_work_items, - int* partition_row_ids, - cudaStream_t builder_stream); +template void launchNodeSplitKernel(const Dataset& dataset, + const NodeWorkItem* work_items, + Split* splits, + const WorkloadInfo* workload_info, + size_t n_blocks_dimx, + size_t n_work_items, + std::int64_t* partition_row_ids, + cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchNodeSplitKernel( - const Dataset& dataset, - const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, - size_t n_blocks_dimx, - size_t n_work_items, - int* partition_row_ids, - cudaStream_t builder_stream); +template void launchNodeSplitKernel(const Dataset& dataset, + const NodeWorkItem* work_items, + Split* splits, + const WorkloadInfo* workload_info, + size_t n_blocks_dimx, + size_t n_work_items, + std::int64_t* partition_row_ids, + cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchNodeSplitKernel( - const Dataset& dataset, - const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, - size_t n_blocks_dimx, - size_t n_work_items, - int* partition_row_ids, - cudaStream_t builder_stream); +template void launchNodeSplitKernel(const Dataset& dataset, + const NodeWorkItem* work_items, + Split* splits, + const WorkloadInfo* workload_info, + size_t n_blocks_dimx, + size_t n_work_items, + std::int64_t* partition_row_ids, + cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchNodeSplitKernel( - const Dataset& dataset, +template void launchNodeSplitKernel( + const Dataset& dataset, const NodeWorkItem* work_items, - Split* splits, - const WorkloadInfo* workload_info, + Split* splits, + const WorkloadInfo* workload_info, size_t n_blocks_dimx, size_t n_work_items, - int* partition_row_ids, + std::int64_t* partition_row_ids, cudaStream_t builder_stream); } // namespace DT diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu index 3a4b3b2fae..b1e7f54969 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = double; using LabelT = double; -using IdxT = int; -using ObjectiveT = RegressionObjectiveFunction; +using ObjectiveT = RegressionObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu index 858c74f515..672066cc9d 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = float; using LabelT = float; -using IdxT = int; -using ObjectiveT = RegressionObjectiveFunction; +using ObjectiveT = RegressionObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu index 063743e7e2..789a5f21ab 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = double; using LabelT = int; -using IdxT = int; -using ObjectiveT = ClassificationObjectiveFunction; +using ObjectiveT = ClassificationObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu index 4e71eeec8a..8ebe17eeaf 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = float; using LabelT = int; -using IdxT = int; -using ObjectiveT = ClassificationObjectiveFunction; +using ObjectiveT = ClassificationObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu index 9741d5b23c..d459c46236 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = double; using LabelT = double; -using IdxT = int; -using ObjectiveT = RegressionObjectiveFunction; +using ObjectiveT = RegressionObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu index c03240265d..21ef9a0c07 100644 --- a/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu +++ b/cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu @@ -11,11 +11,10 @@ namespace ML { namespace DT { using DataT = float; using LabelT = float; -using IdxT = int; -using ObjectiveT = RegressionObjectiveFunction; +using ObjectiveT = RegressionObjectiveFunction; using BinT = typename ObjectiveT::BinT; -using DatasetT = Dataset; -using NodeT = SparseTreeNode; +using DatasetT = Dataset; +using NodeT = SparseTreeNode; // Explicit instantiations are split across separate .cu files to increase compilation parallelism. template void launchLeafKernel( @@ -29,30 +28,30 @@ template void launchLeafKernel( cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchBuildHistogramsKernel( +template void launchBuildHistogramsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, + const Quantiles& quantiles, const NodeWorkItem* work_items, - IdxT colStart, - const IdxT* column_samples, + std::int64_t colStart, + const std::int64_t* column_samples, ObjectiveT& objective, - const WorkloadInfo* workload_info, + const WorkloadInfo* workload_info, dim3 histogram_grid, const SharedMemoryConfig& split_smem_config, cudaStream_t builder_stream); // Explicit instantiations are split across separate .cu files to increase compilation parallelism. -template void launchFindBestSplitsKernel( +template void launchFindBestSplitsKernel( BinT* histograms, - IdxT n_bins, + std::int64_t n_bins, const DatasetT& dataset, - const Quantiles& quantiles, - IdxT colStart, - const IdxT* column_samples, + const Quantiles& quantiles, + std::int64_t colStart, + const std::int64_t* column_samples, int* mutex, - volatile Split* splits, + volatile Split* splits, ObjectiveT& objective, dim3 split_grid, cudaStream_t builder_stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/objectives.cuh b/cpp/src/decisiontree/batched-levelalgo/objectives.cuh index b5d030c89d..4cd6ba69a7 100644 --- a/cpp/src/decisiontree/batched-levelalgo/objectives.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/objectives.cuh @@ -17,32 +17,35 @@ namespace ML { namespace DT { -template +template class ClassificationObjectiveFunction { public: using DataT = DataT_; using LabelT = LabelT_; - using IdxT = IdxT_; using BinT = std::conditional_t; static constexpr bool weighted = weighted_; private: - IdxT nclasses; - IdxT min_samples_leaf; + std::int64_t nclasses; + std::int64_t min_samples_leaf; CRITERION criterion; DataT min_impurity_decrease; - HDI double WeightAt(BinT const* hist, IdxT i, IdxT n_bins) const + HDI double WeightAt(BinT const* hist, std::int64_t i, std::int64_t n_bins) const { double weight = 0.0; - for (IdxT j = 0; j < nclasses; ++j) { + for (std::int64_t j = 0; j < nclasses; ++j) { weight += hist[n_bins * j + i].Weight(); } return weight; } - HDI DataT - GiniGain(BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT GiniGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { constexpr DataT One = DataT(1.0); auto total_weight = WeightAt(hist, n_bins - 1, n_bins); @@ -57,7 +60,7 @@ class ClassificationObjectiveFunction { auto invRight = One / DataT(right_weight); auto gain = DataT(0.0); - for (IdxT j = 0; j < nclasses; ++j) { + for (std::int64_t j = 0; j < nclasses; ++j) { double val_i = 0.0; auto lval_i = hist[n_bins * j + i].Weight(); auto lval = DataT(lval_i); @@ -77,8 +80,12 @@ class ClassificationObjectiveFunction { return gain; } - HDI DataT - EntropyGain(BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT EntropyGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { auto total_weight = WeightAt(hist, n_bins - 1, n_bins); auto left_weight = WeightAt(hist, i, n_bins); @@ -91,7 +98,7 @@ class ClassificationObjectiveFunction { auto invLeft{DataT(1.0) / DataT(left_weight)}; auto invRight{DataT(1.0) / DataT(right_weight)}; auto invLen{DataT(1.0) / DataT(total_weight)}; - for (IdxT c = 0; c < nclasses; ++c) { + for (std::int64_t c = 0; c < nclasses; ++c) { double val_i = 0.0; auto lval_i = hist[n_bins * c + i].Weight(); if (lval_i != 0) { @@ -119,8 +126,8 @@ class ClassificationObjectiveFunction { public: HDI DataT GainPerSplit(BinT const* hist, - IdxT i, - IdxT n_bins, + std::int64_t i, + std::int64_t n_bins, std::int64_t len, std::int64_t nLeft, std::int64_t nRight) const @@ -136,8 +143,8 @@ class ClassificationObjectiveFunction { } } - HDI ClassificationObjectiveFunction(IdxT nclasses, - IdxT min_samples_leaf, + HDI ClassificationObjectiveFunction(std::int64_t nclasses, + std::int64_t min_samples_leaf, CRITERION criterion, DataT min_impurity_decrease = DataT{0}) : nclasses(nclasses), @@ -147,11 +154,15 @@ class ClassificationObjectiveFunction { { } - DI IdxT NumClasses() const { return nclasses; } + DI std::int64_t NumClasses() const { return nclasses; } template - DI void IncrementHistogram( - BinT* histogram, IdxT n_bins, IdxT bin, LabelT label, const DatasetT& dataset, IdxT row) const + DI void IncrementHistogram(BinT* histogram, + int n_bins, + int bin, + LabelT label, + const DatasetT& dataset, + std::int64_t row) const { double weight = 1.0; if constexpr (weighted) { @@ -160,11 +171,14 @@ class ClassificationObjectiveFunction { BinT::IncrementHistogram(histogram, n_bins, bin, label, weight); } - DI Split Gain( - BinT const* shist, DataT const* squantiles, IdxT col, std::int64_t len, IdxT n_bins) const + DI Split Gain(BinT const* shist, + DataT const* squantiles, + std::int64_t col, + std::int64_t len, + std::int64_t n_bins) const { - Split sp; - for (IdxT i = threadIdx.x; i < n_bins; i += blockDim.x) { + Split sp; + for (std::int64_t i = threadIdx.x; i < n_bins; i += blockDim.x) { auto nLeft = detail::CountLeft(shist, i, n_bins, nclasses); auto nRight = len - nLeft; if (nLeft >= static_cast(min_samples_leaf) && @@ -195,23 +209,26 @@ class ClassificationObjectiveFunction { } }; -template +template class RegressionObjectiveFunction { public: using DataT = DataT_; using LabelT = LabelT_; - using IdxT = IdxT_; using BinT = std::conditional_t; static constexpr bool weighted = weighted_; private: - IdxT min_samples_leaf; + std::int64_t min_samples_leaf; CRITERION criterion; DataT min_impurity_decrease; static constexpr auto eps_ = 10 * std::numeric_limits::epsilon(); - HDI DataT - MSEGain(BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT MSEGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { auto parent_weight = hist[n_bins - 1].Weight(); auto left_weight = hist[i].Weight(); @@ -233,8 +250,12 @@ class RegressionObjectiveFunction { return gain; } - HDI DataT - PoissonGain(BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT PoissonGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { auto parent_weight = hist[n_bins - 1].Weight(); auto left_weight = hist[i].Weight(); @@ -261,8 +282,12 @@ class RegressionObjectiveFunction { return gain; } - HDI DataT - GammaGain(BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT GammaGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { auto parent_weight = hist[n_bins - 1].Weight(); auto left_weight = hist[i].Weight(); @@ -289,8 +314,12 @@ class RegressionObjectiveFunction { return gain; } - HDI DataT InverseGaussianGain( - BinT const* hist, IdxT i, IdxT n_bins, std::int64_t, std::int64_t, std::int64_t) const + HDI DataT InverseGaussianGain(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t, + std::int64_t, + std::int64_t) const { auto parent_weight = hist[n_bins - 1].Weight(); auto left_weight = hist[i].Weight(); @@ -318,8 +347,8 @@ class RegressionObjectiveFunction { public: HDI DataT GainPerSplit(BinT const* hist, - IdxT i, - IdxT n_bins, + std::int64_t i, + std::int64_t n_bins, std::int64_t len, std::int64_t nLeft, std::int64_t nRight) const @@ -338,8 +367,8 @@ class RegressionObjectiveFunction { } } - HDI RegressionObjectiveFunction(IdxT, - IdxT min_samples_leaf, + HDI RegressionObjectiveFunction(std::int64_t, + std::int64_t min_samples_leaf, CRITERION criterion, DataT min_impurity_decrease = DataT{0}) : min_samples_leaf(min_samples_leaf), @@ -348,11 +377,15 @@ class RegressionObjectiveFunction { { } - DI IdxT NumClasses() const { return 1; } + DI std::int64_t NumClasses() const { return 1; } template - DI void IncrementHistogram( - BinT* histogram, IdxT n_bins, IdxT bin, LabelT label, const DatasetT& dataset, IdxT row) const + DI void IncrementHistogram(BinT* histogram, + int n_bins, + int bin, + LabelT label, + const DatasetT& dataset, + std::int64_t row) const { double weight = 1.0; if constexpr (weighted) { @@ -361,12 +394,15 @@ class RegressionObjectiveFunction { BinT::IncrementHistogram(histogram, n_bins, bin, label, weight); } - DI Split Gain( - BinT const* shist, DataT const* squantiles, IdxT col, std::int64_t len, IdxT n_bins) const + DI Split Gain(BinT const* shist, + DataT const* squantiles, + std::int64_t col, + std::int64_t len, + std::int64_t n_bins) const { - Split sp; - for (IdxT i = threadIdx.x; i < n_bins; i += blockDim.x) { - auto nLeft = detail::CountLeft(shist, i, n_bins, IdxT{1}); + Split sp; + for (std::int64_t i = threadIdx.x; i < n_bins; i += blockDim.x) { + auto nLeft = detail::CountLeft(shist, i, n_bins, std::int64_t{1}); auto nRight = len - nLeft; if (nLeft >= static_cast(min_samples_leaf) && nRight >= static_cast(min_samples_leaf)) { diff --git a/cpp/src/decisiontree/batched-levelalgo/quantiles.cuh b/cpp/src/decisiontree/batched-levelalgo/quantiles.cuh index 93b2d60177..f28aefbf18 100644 --- a/cpp/src/decisiontree/batched-levelalgo/quantiles.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/quantiles.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -115,8 +115,8 @@ struct QuantileResult { rmm::device_uvector quantiles_array; rmm::device_uvector n_bins_array; - Quantiles view() & { return {quantiles_array.data(), n_bins_array.data()}; } - Quantiles view() && = delete; + Quantiles view() & { return {quantiles_array.data(), n_bins_array.data()}; } + Quantiles view() && = delete; }; /** @@ -197,12 +197,10 @@ CUML_EXPORT QuantileResult computeQuantiles(const raft::handle_t& handle, rmm::device_uvector sampled_columns(total_sample_values, stream); rmm::device_uvector sorted_samples(total_sample_values, stream); - int n_threads = 256; - auto segment_offsets = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [sample_count] __host__ __device__(std::int64_t col) { - return col * static_cast(sample_count); - }); + int n_threads = 256; + auto segment_offsets = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + [sample_count] __host__ __device__(std::int64_t col) { return col * sample_count; }); rmm::device_uvector quantiles_array(ML::checked_mul(n_cols, max_n_bins), stream); rmm::device_uvector n_bins_array(n_cols, stream); diff --git a/cpp/src/decisiontree/batched-levelalgo/quantiles.h b/cpp/src/decisiontree/batched-levelalgo/quantiles.h index abfcec3ab8..3b65de1a3b 100644 --- a/cpp/src/decisiontree/batched-levelalgo/quantiles.h +++ b/cpp/src/decisiontree/batched-levelalgo/quantiles.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2022, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -8,12 +8,12 @@ namespace ML { namespace DT { -template +template struct Quantiles { /** quantiles computed for each feature of dataset in col-major */ DataT* quantiles_array; /** The number of bins used for quantiles of each feature*/ - IdxT* n_bins_array; + int* n_bins_array; }; } // namespace DT diff --git a/cpp/src/decisiontree/batched-levelalgo/split.cuh b/cpp/src/decisiontree/batched-levelalgo/split.cuh index 1c040246f9..5ff7ac356e 100644 --- a/cpp/src/decisiontree/batched-levelalgo/split.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/split.cuh @@ -16,11 +16,14 @@ namespace ML { namespace DT { namespace detail { -template -DI std::int64_t CountLeft(BinT const* hist, IdxT i, IdxT n_bins, IdxT n_outputs) +template +DI std::int64_t CountLeft(BinT const* hist, + std::int64_t i, + std::int64_t n_bins, + std::int64_t n_outputs) { auto nLeft = hist[i].Count(); - for (IdxT j = 1; j < n_outputs; ++j) { + for (std::int64_t j = 1; j < n_outputs; ++j) { nLeft += hist[n_bins * j + i].Count(); } return static_cast(nLeft); @@ -33,16 +36,16 @@ DI std::int64_t CountLeft(BinT const* hist, IdxT i, IdxT n_bins, IdxT n_outputs) * * @tparam DataT input data type */ -template +template struct Split { - typedef Split SplitT; + typedef Split SplitT; /** start with this as the initial gain */ static constexpr DataT Min = -std::numeric_limits::max(); /** threshold to compare in this node */ DataT quesval; /** feature index */ - IdxT colid; + std::int64_t colid; /** best info gain on this node */ DataT best_metric_val; /** global number of samples in the left child */ @@ -50,9 +53,9 @@ struct Split { /** rank-local number of samples in the left child */ std::int64_t local_nLeft; /** first quantile index in an inclusive range of training-equivalent splits */ - IdxT split_start; + std::int64_t split_start; /** last quantile index in an inclusive range of training-equivalent splits */ - IdxT split_end; + std::int64_t split_end; HDI Split() { @@ -76,27 +79,27 @@ struct Split { return *this; } - HDI bool IsValid() const { return colid != static_cast(-1); } + HDI bool IsValid() const { return colid != static_cast(-1); } DI bool has_valid_split_range() const { - return split_start >= IdxT{0} && split_end >= split_start; + return split_start >= std::int64_t{0} && split_end >= split_start; } DI bool can_merge_equivalent_split_range(std::int64_t other_global_nLeft, - IdxT other_split_start, - IdxT other_split_end) const + std::int64_t other_split_start, + std::int64_t other_split_end) const { return global_nLeft == other_global_nLeft && has_valid_split_range() && - other_split_start >= IdxT{0} && other_split_end >= other_split_start; + other_split_start >= std::int64_t{0} && other_split_end >= other_split_start; } // Extend the candidate's inclusive range of training-equivalent split // thresholds. `quesval` tracks the upper edge only to preserve the existing // threshold tie-break against candidates outside this equivalent range. DI void merge_equivalent_split_range(DataT other_quesval, - IdxT other_split_start, - IdxT other_split_end) + std::int64_t other_split_start, + std::int64_t other_split_end) { split_start = other_split_start < split_start ? other_split_start : split_start; split_end = other_split_end > split_end ? other_split_end : split_end; @@ -104,11 +107,11 @@ struct Split { } DI bool replace_with(DataT other_quesval, - IdxT other_colid, + std::int64_t other_colid, DataT other_best_metric_val, std::int64_t other_global_nLeft, - IdxT other_split_start, - IdxT other_split_end) + std::int64_t other_split_start, + std::int64_t other_split_end) { quesval = other_quesval; colid = other_colid; @@ -123,7 +126,7 @@ struct Split { // Several thresholds can be equally good for the training data while still // routing future inference values differently. Select the middle split in // that equivalent range so deterministic tie-breaking does not pick an edge. - DI void select_split_range_midpoint(DataT const* quantiles, IdxT n_bins) + DI void select_split_range_midpoint(DataT const* quantiles, std::int64_t n_bins) { if (has_valid_split_range() && split_end < n_bins) { auto bin = split_start + (split_end - split_start + 1) / 2; @@ -140,11 +143,11 @@ struct Split { * on global counts, and local counts are filled just before partitioning. */ DI bool update(DataT other_quesval, - IdxT other_colid, + std::int64_t other_colid, DataT other_best_metric_val, std::int64_t other_global_nLeft, - IdxT other_split_start, - IdxT other_split_end) + std::int64_t other_split_start, + std::int64_t other_split_end) { // Primary ordering: higher gain wins; lower or unordered gain loses. if (other_best_metric_val > best_metric_val) { @@ -191,10 +194,10 @@ struct Split { } DI bool update(DataT other_quesval, - IdxT other_colid, + std::int64_t other_colid, DataT other_best_metric_val, std::int64_t other_global_nLeft, - IdxT other_bin) + std::int64_t other_bin) { return update( other_quesval, other_colid, other_best_metric_val, other_global_nLeft, other_bin, other_bin); @@ -229,8 +232,11 @@ struct Split { * @note all threads in the block must enter this function together. At the * end thread0 will contain the best split. */ - DI void evalBestSplit( - SplitT* split_scratch, volatile SplitT* split, int* mutex, DataT const* quantiles, IdxT n_bins) + DI void evalBestSplit(SplitT* split_scratch, + volatile SplitT* split, + int* mutex, + DataT const* quantiles, + std::int64_t n_bins) { warpReduce(); auto warp = threadIdx.x / raft::WarpSize; @@ -281,17 +287,17 @@ struct Split { * @param[in] len length of this array * @param[in] s cuda stream where to schedule work */ -template -void initSplit(Split* splits, IdxT len, cudaStream_t s) +template +void initSplit(Split* splits, std::int64_t len, cudaStream_t s) { - auto op = [] __device__(Split * ptr, IdxT idx) { *ptr = Split(); }; - raft::linalg::writeOnlyUnaryOp, decltype(op), IdxT, TPB>(splits, len, op, s); + auto op = [] __device__(Split * ptr, std::int64_t idx) { *ptr = Split(); }; + raft::linalg::writeOnlyUnaryOp, decltype(op), std::int64_t, TPB>(splits, len, op, s); } -template -void printSplits(Split* splits, IdxT len, cudaStream_t s) +template +void printSplits(Split* splits, std::int64_t len, cudaStream_t s) { - auto op = [] __device__(Split * ptr, IdxT idx) { + auto op = [] __device__(Split * ptr, std::int64_t idx) { printf( "quesval = %e, colid = %lld, best_metric_val = %e, global_nLeft = %lld, " "local_nLeft = %lld, split_range = [%lld, %lld]\n", @@ -303,7 +309,7 @@ void printSplits(Split* splits, IdxT len, cudaStream_t s) static_cast(ptr->split_start), static_cast(ptr->split_end)); }; - raft::linalg::writeOnlyUnaryOp, decltype(op), IdxT, TPB>(splits, len, op, s); + raft::linalg::writeOnlyUnaryOp, decltype(op), std::int64_t, TPB>(splits, len, op, s); RAFT_CUDA_TRY(cudaDeviceSynchronize()); } diff --git a/cpp/src/decisiontree/decisiontree.cuh b/cpp/src/decisiontree/decisiontree.cuh index 80826959a4..ddc0c72594 100644 --- a/cpp/src/decisiontree/decisiontree.cuh +++ b/cpp/src/decisiontree/decisiontree.cuh @@ -21,6 +21,7 @@ #include #include +#include #include #include #include @@ -87,7 +88,7 @@ std::string get_node_text(const std::string& prefix, int idx, bool isLeft) { - const SparseTreeNode& node = tree->sparsetree[idx]; + const auto& node = tree->sparsetree[idx]; std::ostringstream oss; @@ -124,7 +125,7 @@ std::string get_node_text(const std::string& prefix, template std::string get_node_json(const std::string& prefix, const TreeMetaDataNode* tree, int idx) { - const SparseTreeNode& node = tree->sparsetree[idx]; + const auto& node = tree->sparsetree[idx]; std::ostringstream oss; if (!node.IsLeaf()) { @@ -181,9 +182,9 @@ tl::Tree build_treelite_tree(const DT::TreeMetaDataNode& rf_tree, next_level_queue.resize(std::max(2 * cur_level_size, next_level_queue.size())); for (size_t i = 0; i < cur_level_size; ++i) { - auto cuml_node_id = cur_level_queue[cur_front].first; - const SparseTreeNode& q_node = rf_tree.sparsetree[cuml_node_id]; - auto tl_node_id = cur_level_queue[cur_front].second; + auto cuml_node_id = cur_level_queue[cur_front].first; + const auto& q_node = rf_tree.sparsetree[cuml_node_id]; + auto tl_node_id = cur_level_queue[cur_front].second; ++cur_front; if (!q_node.IsLeaf()) { @@ -236,15 +237,15 @@ class DecisionTree { const raft::handle_t& handle, const cudaStream_t s, const DataT* data, - const int ncols, - const int nrows, + const std::int64_t ncols, + const std::int64_t nrows, const LabelT* labels, - rmm::device_uvector* row_ids, + rmm::device_uvector* row_ids, int unique_labels, DecisionTreeParams params, uint64_t seed, - const Quantiles& quantiles, - int treeid, + const Quantiles& quantiles, + std::int64_t treeid, const double* sample_weight = nullptr, bool row_major = false) { @@ -254,41 +255,40 @@ class DecisionTree { (std::numeric_limits::is_integer) ? CRITERION::GINI : CRITERION::MSE; params.split_criterion = default_criterion; } - using IdxT = int; // Dispatch objective family. The objective object switches on the criterion at runtime. if (not std::is_same::value and (params.split_criterion == CRITERION::GINI || params.split_criterion == CRITERION::ENTROPY)) { if (sample_weight != nullptr) { - return Builder>(handle, - s, - treeid, - seed, - params, - data, - labels, - sample_weight, - nrows, - ncols, - row_ids, - unique_labels, - quantiles, - row_major) + return Builder>(handle, + s, + treeid, + seed, + params, + data, + labels, + sample_weight, + nrows, + ncols, + row_ids, + unique_labels, + quantiles, + row_major) .train(); } - return Builder>(handle, - s, - treeid, - seed, - params, - data, - labels, - sample_weight, - nrows, - ncols, - row_ids, - unique_labels, - quantiles, - row_major) + return Builder>(handle, + s, + treeid, + seed, + params, + data, + labels, + sample_weight, + nrows, + ncols, + row_ids, + unique_labels, + quantiles, + row_major) .train(); } else if (std::is_same::value and (params.split_criterion == CRITERION::MSE || @@ -296,36 +296,36 @@ class DecisionTree { params.split_criterion == CRITERION::GAMMA || params.split_criterion == CRITERION::INVERSE_GAUSSIAN)) { if (sample_weight != nullptr) { - return Builder>(handle, - s, - treeid, - seed, - params, - data, - labels, - sample_weight, - nrows, - ncols, - row_ids, - unique_labels, - quantiles, - row_major) + return Builder>(handle, + s, + treeid, + seed, + params, + data, + labels, + sample_weight, + nrows, + ncols, + row_ids, + unique_labels, + quantiles, + row_major) .train(); } - return Builder>(handle, - s, - treeid, - seed, - params, - data, - labels, - sample_weight, - nrows, - ncols, - row_ids, - unique_labels, - quantiles, - row_major) + return Builder>(handle, + s, + treeid, + seed, + params, + data, + labels, + sample_weight, + nrows, + ncols, + row_ids, + unique_labels, + quantiles, + row_major) .train(); } else { ASSERT(false, "Unknown split criterion."); diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh index e9ee8aa3bd..8ce5992eef 100644 --- a/cpp/src/randomforest/randomforest.cuh +++ b/cpp/src/randomforest/randomforest.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -42,6 +42,7 @@ #define omp_get_max_threads() 1 #endif +#include #include #include @@ -64,8 +65,8 @@ class RowSampler { public: RowSampler(const raft::handle_t& handle, const RF_params& rf_params, - int n_rows, - int n_sampled_rows, + std::int64_t n_rows, + std::int64_t n_sampled_rows, int n_streams, bool* bootstrap_masks, const double* sample_weight) @@ -82,7 +83,7 @@ class RowSampler { "bootstrap_masks must be a GPU pointer"); validate_sample_weight(handle, sample_weight_, n_rows_); if (use_weighted_bootstrap()) { - sample_weight_cdf_.resize(n_rows_, handle.get_stream()); + sample_weight_cdf_.resize(ML::narrow_cast(n_rows_), handle.get_stream()); thrust::inclusive_scan(rmm::exec_policy(handle.get_stream()), sample_weight_, sample_weight_ + n_rows_, @@ -95,11 +96,12 @@ class RowSampler { "sample_weight values must contain at least one positive value"); } // Use a deque instead of vector because device_uvector has a deleted copy constructor. + auto const n_sampled_rows_size = ML::narrow_cast(n_sampled_rows_); for (int i = 0; i < n_streams; i++) { auto stream = handle.get_stream_from_stream_pool(i); - selected_rows_.emplace_back(n_sampled_rows_, stream); + selected_rows_.emplace_back(n_sampled_rows_size, stream); if (use_weighted_bootstrap()) { - weighted_draw_scratch_.emplace_back(n_sampled_rows_, stream); + weighted_draw_scratch_.emplace_back(n_sampled_rows_size, stream); } } } @@ -107,7 +109,7 @@ class RowSampler { RowSampler(const RowSampler&) = delete; RowSampler& operator=(const RowSampler&) = delete; - rmm::device_uvector& sample(int tree_id, int stream_id, cudaStream_t stream) + rmm::device_uvector& sample(int tree_id, int stream_id, cudaStream_t stream) { raft::common::nvtx::range fun_scope("bootstrapping row IDs @randomforest.cuh"); @@ -139,12 +141,12 @@ class RowSampler { selected_rows.begin()); } else if (bootstrap_) { // Draw bootstrap rows uniformly when there are no sample weights. - raft::random::uniformInt( + raft::random::uniformInt( stream_resources, rng_state, selected_rows.data(), selected_rows.size(), 0, n_rows_); } else if (sample_weight_ != nullptr) { // Remove zero-weight rows from the non-bootstrap row set. - selected_rows.resize(n_sampled_rows_, stream); - auto rows_begin = thrust::make_counting_iterator(0); + selected_rows.resize(ML::narrow_cast(n_sampled_rows_), stream); + auto rows_begin = thrust::make_counting_iterator(0); auto selected_rows_end = thrust::copy_if(rmm::exec_policy(stream), rows_begin, rows_begin + n_rows_, @@ -155,7 +157,7 @@ class RowSampler { ASSERT(n_selected > 0, "sample_weight values must contain at least one positive value"); selected_rows.resize(n_selected, stream); } else { - selected_rows.resize(n_sampled_rows_, stream); + selected_rows.resize(ML::narrow_cast(n_sampled_rows_), stream); thrust::sequence(rmm::exec_policy(stream), selected_rows.begin(), selected_rows.end()); } @@ -168,7 +170,7 @@ class RowSampler { private: void store_bootstrap_mask(int tree_id, - rmm::device_uvector& selected_rows, + rmm::device_uvector& selected_rows, cudaStream_t stream) { if (bootstrap_masks_ == nullptr) { return; } @@ -198,7 +200,7 @@ class RowSampler { static void validate_sample_weight(const raft::handle_t& handle, const double* sample_weight, - int n_rows) + std::int64_t n_rows) { ASSERT(sample_weight == nullptr || DT::is_dev_ptr(sample_weight), "sample_weight must be a GPU pointer"); @@ -215,13 +217,13 @@ class RowSampler { bool bootstrap_; uint64_t seed_; - int n_rows_; - int n_sampled_rows_; + std::int64_t n_rows_; + std::int64_t n_sampled_rows_; bool* bootstrap_masks_; const double* sample_weight_; double sample_weight_sum_; rmm::device_uvector sample_weight_cdf_; - std::deque> selected_rows_; + std::deque> selected_rows_; std::deque> weighted_draw_scratch_; }; } // namespace detail @@ -296,10 +298,12 @@ class RandomForest { { raft::common::nvtx::range fun_scope("RandomForest::fit @randomforest.cuh"); this->error_checking(input, labels, n_rows, n_cols, false); - const raft::handle_t& handle = user_handle; - int n_sampled_rows = 0; + const raft::handle_t& handle = user_handle; + std::int64_t const n_rows_i64 = n_rows; + std::int64_t n_sampled_rows = 0; if (this->rf_params.bootstrap) { - n_sampled_rows = std::round(this->rf_params.max_samples * n_rows); + n_sampled_rows = + static_cast(std::round(this->rf_params.max_samples * n_rows_i64)); } else { if (this->rf_params.max_samples != 1.0) { CUML_LOG_WARN( @@ -307,7 +311,7 @@ class RandomForest { "whole dataset is used for building each tree"); this->rf_params.max_samples = 1.0; } - n_sampled_rows = n_rows; + n_sampled_rows = n_rows_i64; } int n_streams = this->rf_params.n_streams; ASSERT(static_cast(n_streams) <= handle.get_stream_pool_size(), @@ -328,8 +332,13 @@ class RandomForest { // n_streams should not be less than n_trees if (this->rf_params.n_trees < n_streams) n_streams = this->rf_params.n_trees; - detail::RowSampler row_sampler( - handle, this->rf_params, n_rows, n_sampled_rows, n_streams, bootstrap_masks, sample_weight); + detail::RowSampler row_sampler(handle, + this->rf_params, + n_rows_i64, + n_sampled_rows, + n_streams, + bootstrap_masks, + sample_weight); forest->n_features = n_cols; diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu index f26c2ea84d..12ac6f8ccb 100644 --- a/cpp/tests/sg/rf_test.cu +++ b/cpp/tests/sg/rf_test.cu @@ -2,6 +2,7 @@ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include #include @@ -1123,57 +1124,6 @@ struct QuantileTestParameters { uint64_t seed; }; -template -class RFQuantileBinsLowerBoundTest : public ::testing::TestWithParam { - public: - void SetUp() override - { - auto params = ::testing::TestWithParam::GetParam(); - - thrust::device_vector data(params.n_rows); - thrust::host_vector h_data(params.n_rows); - thrust::host_vector h_quantiles(params.max_n_bins); - raft::random::Rng r(8); - r.normal(data.data().get(), data.size(), T(0.0), T(2.0), nullptr); - auto stream_pool = std::make_shared(1); - raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); - - // computing the quantiles - auto quantile_result = - DT::computeQuantiles(handle, data.data().get(), params.max_n_bins, params.n_rows, 1); - auto quantiles = quantile_result.view(); - - raft::update_host( - h_quantiles.data(), quantiles.quantiles_array, params.max_n_bins, handle.get_stream()); - - int n_unique_bins; - raft::copy(&n_unique_bins, quantiles.n_bins_array, 1, handle.get_stream()); - if (n_unique_bins < params.max_n_bins) { - return; // almost impossible that this happens, skip if so - } - - h_data = data; - for (std::size_t i = 0; i < h_data.size(); ++i) { - auto d = h_data[i]; - // golden lower bound from thrust - auto golden_lb = - thrust::lower_bound( - thrust::seq, h_quantiles.data(), h_quantiles.data() + params.max_n_bins, d) - - h_quantiles.data(); - // lower bound from custom lower_bound impl - auto lb = DT::lower_bound(h_quantiles.data(), params.max_n_bins, d); - if (golden_lb == params.max_n_bins) { - ASSERT_EQ(lb, params.max_n_bins - 1) - << "custom lower_bound should clamp values above the last quantile to the last bin" - << std::endl; - continue; - } - ASSERT_EQ(golden_lb, lb) - << "custom lower_bound method is inconsistent with thrust::lower_bound" << std::endl; - } - } -}; - template class RFQuantileTest : public ::testing::TestWithParam { public: @@ -1362,21 +1312,21 @@ class RFSampledQuantileDeterminismTest : public ::testing::TestWithParam +template __global__ void objectiveGainKernel(BinT const* hist, DataT const* quantiles, - DT::Split* out, + DT::Split* out, int* mutex, ObjectiveT objective, - IdxT col, - IdxT len, - IdxT n_bins) + std::int64_t col, + std::int64_t len, + std::int64_t n_bins) { - __shared__ __align__(alignof( - DT::Split)) unsigned char split_scratch_storage[sizeof(DT::Split)]; - auto* split_scratch = reinterpret_cast*>(split_scratch_storage); + __shared__ __align__( + alignof(DT::Split)) unsigned char split_scratch_storage[sizeof(DT::Split)]; + auto* split_scratch = reinterpret_cast*>(split_scratch_storage); if (threadIdx.x == 0) { - *out = DT::Split(); + *out = DT::Split(); *mutex = 0; } __syncthreads(); @@ -1388,10 +1338,9 @@ __global__ void objectiveGainKernel(BinT const* hist, TEST(RFEquivalentSplitRangeTest, ClassificationChoosesUpperMiddleBin) { - using DataT = float; - using IdxT = int; - constexpr IdxT len = 8; - constexpr IdxT n_bins = 6; + using DataT = float; + constexpr std::int64_t len = 8; + constexpr std::int64_t n_bins = 6; auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); @@ -1414,21 +1363,21 @@ TEST(RFEquivalentSplitRangeTest, ClassificationChoosesUpperMiddleBin) thrust::device_vector hist(h_hist.begin(), h_hist.end()); thrust::device_vector quantiles(h_quantiles.begin(), h_quantiles.end()); - thrust::device_vector> split(1); + thrust::device_vector> split(1); thrust::device_vector mutex(1); - DT::ClassificationObjectiveFunction objective(2, 1, CRITERION::GINI); + DT::ClassificationObjectiveFunction objective(2, 1, CRITERION::GINI); objectiveGainKernel<<<1, 32, 0, handle.get_stream()>>>(hist.data().get(), quantiles.data().get(), split.data().get(), mutex.data().get(), objective, - IdxT{0}, + std::int64_t{0}, len, n_bins); RAFT_CUDA_TRY(cudaGetLastError()); - DT::Split h_split; + DT::Split h_split; RAFT_CUDA_TRY(cudaMemcpyAsync( &h_split, split.data().get(), sizeof(h_split), cudaMemcpyDeviceToHost, handle.get_stream())); handle.sync_stream(); @@ -1442,10 +1391,9 @@ TEST(RFEquivalentSplitRangeTest, ClassificationChoosesUpperMiddleBin) TEST(RFEquivalentSplitRangeTest, RegressionChoosesUpperMiddleBin) { - using DataT = float; - using IdxT = int; - constexpr IdxT len = 8; - constexpr IdxT n_bins = 6; + using DataT = float; + constexpr std::int64_t len = 8; + constexpr std::int64_t n_bins = 6; auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); @@ -1462,21 +1410,21 @@ TEST(RFEquivalentSplitRangeTest, RegressionChoosesUpperMiddleBin) thrust::device_vector hist(h_hist.begin(), h_hist.end()); thrust::device_vector quantiles(h_quantiles.begin(), h_quantiles.end()); - thrust::device_vector> split(1); + thrust::device_vector> split(1); thrust::device_vector mutex(1); - DT::RegressionObjectiveFunction objective(1, 1, CRITERION::MSE); + DT::RegressionObjectiveFunction objective(1, 1, CRITERION::MSE); objectiveGainKernel<<<1, 32, 0, handle.get_stream()>>>(hist.data().get(), quantiles.data().get(), split.data().get(), mutex.data().get(), objective, - IdxT{0}, + std::int64_t{0}, len, n_bins); RAFT_CUDA_TRY(cudaGetLastError()); - DT::Split h_split; + DT::Split h_split; RAFT_CUDA_TRY(cudaMemcpyAsync( &h_split, split.data().get(), sizeof(h_split), cudaMemcpyDeviceToHost, handle.get_stream())); handle.sync_stream(); @@ -1581,16 +1529,6 @@ typedef RFQuantileTest RFQuantileTestD; TEST_P(RFQuantileTestD, test) {} INSTANTIATE_TEST_CASE_P(RfTests, RFQuantileTestD, ::testing::ValuesIn(inputs)); -// float type quantile bins lower bounds test -typedef RFQuantileBinsLowerBoundTest RFQuantileBinsLowerBoundTestF; -TEST_P(RFQuantileBinsLowerBoundTestF, test) {} -INSTANTIATE_TEST_CASE_P(RfTests, RFQuantileBinsLowerBoundTestF, ::testing::ValuesIn(inputs)); - -// double type quantile bins lower bounds test -typedef RFQuantileBinsLowerBoundTest RFQuantileBinsLowerBoundTestD; -TEST_P(RFQuantileBinsLowerBoundTestD, test) {} -INSTANTIATE_TEST_CASE_P(RfTests, RFQuantileBinsLowerBoundTestD, ::testing::ValuesIn(inputs)); - // float type quantile variable binning test typedef RFQuantileVariableBinsTest RFQuantileVariableBinsTestF; TEST_P(RFQuantileVariableBinsTestF, test) {} @@ -1949,7 +1887,6 @@ class ObjectiveTest : public ::testing::TestWithParam { using ObjectiveT = typename ObjectiveConfig::ObjectiveT; typedef typename ObjectiveT::DataT DataT; typedef typename ObjectiveT::LabelT LabelT; - typedef typename ObjectiveT::IdxT IdxT; typedef typename ObjectiveT::BinT BinT; static constexpr auto eps_ = 10 * std::numeric_limits::epsilon(); @@ -2387,11 +2324,11 @@ class ObjectiveTest : public ::testing::TestWithParam { return DataT(0.0); } - auto NumLeftOfBin(std::vector const& cdf_hist, IdxT idx) + auto NumLeftOfBin(std::vector const& cdf_hist, std::int64_t idx) { - auto count{IdxT(0)}; + auto count{std::int64_t(0)}; for (auto c = 0; c < params.n_classes; ++c) { - count += static_cast(cdf_hist[params.max_n_bins * c + idx].Count()); + count += static_cast(cdf_hist[params.max_n_bins * c + idx].Count()); } return count; } @@ -2423,7 +2360,7 @@ class ObjectiveTest : public ::testing::TestWithParam { TEST(WeightedObjectiveEdgeCases, ClassificationRejectsZeroWeightChild) { - using ObjectiveT = ClassificationObjectiveFunction; + using ObjectiveT = ClassificationObjectiveFunction; WeightedClassificationBin hist[]{{1, 0.0}, {1, 0.0}, {0, 0.0}, {1, 1.0}}; CRITERION criteria[] = {CRITERION::GINI, CRITERION::ENTROPY}; @@ -2436,7 +2373,7 @@ TEST(WeightedObjectiveEdgeCases, ClassificationRejectsZeroWeightChild) TEST(WeightedObjectiveEdgeCases, RegressionRejectsZeroWeightChild) { - using ObjectiveT = RegressionObjectiveFunction; + using ObjectiveT = RegressionObjectiveFunction; WeightedRegressionBin hist[]{{0.0, 1, 0.0}, {2.0, 2, 1.0}}; CRITERION criteria[] = { CRITERION::MSE, CRITERION::POISSON, CRITERION::GAMMA, CRITERION::INVERSE_GAUSSIAN}; @@ -2498,28 +2435,28 @@ const std::vector gini_objective_test_parameters = { // mse objective test typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::MSE>> + ObjectiveTestConfig, CRITERION::MSE>> MSEObjectiveTestD; TEST_P(MSEObjectiveTestD, MSEObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, MSEObjectiveTestD, ::testing::ValuesIn(mse_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::MSE>> + ObjectiveTestConfig, CRITERION::MSE>> MSEObjectiveTestF; TEST_P(MSEObjectiveTestF, MSEObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, MSEObjectiveTestF, ::testing::ValuesIn(mse_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::MSE>> + ObjectiveTestConfig, CRITERION::MSE>> WeightedMSEObjectiveTestD; TEST_P(WeightedMSEObjectiveTestD, MSEObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedMSEObjectiveTestD, ::testing::ValuesIn(mse_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::MSE>> + ObjectiveTestConfig, CRITERION::MSE>> WeightedMSEObjectiveTestF; TEST_P(WeightedMSEObjectiveTestF, MSEObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2528,28 +2465,28 @@ INSTANTIATE_TEST_CASE_P(RfTests, // poisson objective test typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::POISSON>> + ObjectiveTestConfig, CRITERION::POISSON>> PoissonObjectiveTestD; TEST_P(PoissonObjectiveTestD, poissonObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, PoissonObjectiveTestD, ::testing::ValuesIn(poisson_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::POISSON>> + ObjectiveTestConfig, CRITERION::POISSON>> PoissonObjectiveTestF; TEST_P(PoissonObjectiveTestF, poissonObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, PoissonObjectiveTestF, ::testing::ValuesIn(poisson_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::POISSON>> + ObjectiveTestConfig, CRITERION::POISSON>> WeightedPoissonObjectiveTestD; TEST_P(WeightedPoissonObjectiveTestD, poissonObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedPoissonObjectiveTestD, ::testing::ValuesIn(poisson_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::POISSON>> + ObjectiveTestConfig, CRITERION::POISSON>> WeightedPoissonObjectiveTestF; TEST_P(WeightedPoissonObjectiveTestF, poissonObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2558,28 +2495,28 @@ INSTANTIATE_TEST_CASE_P(RfTests, // gamma objective test typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GAMMA>> + ObjectiveTestConfig, CRITERION::GAMMA>> GammaObjectiveTestD; TEST_P(GammaObjectiveTestD, GammaObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, GammaObjectiveTestD, ::testing::ValuesIn(gamma_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GAMMA>> + ObjectiveTestConfig, CRITERION::GAMMA>> GammaObjectiveTestF; TEST_P(GammaObjectiveTestF, GammaObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, GammaObjectiveTestF, ::testing::ValuesIn(gamma_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GAMMA>> + ObjectiveTestConfig, CRITERION::GAMMA>> WeightedGammaObjectiveTestD; TEST_P(WeightedGammaObjectiveTestD, GammaObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedGammaObjectiveTestD, ::testing::ValuesIn(gamma_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GAMMA>> + ObjectiveTestConfig, CRITERION::GAMMA>> WeightedGammaObjectiveTestF; TEST_P(WeightedGammaObjectiveTestF, GammaObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2587,29 +2524,29 @@ INSTANTIATE_TEST_CASE_P(RfTests, ::testing::ValuesIn(gamma_objective_test_parameters)); // InvGauss objective test -typedef ObjectiveTest, - CRITERION::INVERSE_GAUSSIAN>> +typedef ObjectiveTest< + ObjectiveTestConfig, CRITERION::INVERSE_GAUSSIAN>> InverseGaussianObjectiveTestD; TEST_P(InverseGaussianObjectiveTestD, InverseGaussianObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, InverseGaussianObjectiveTestD, ::testing::ValuesIn(invgauss_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::INVERSE_GAUSSIAN>> + ObjectiveTestConfig, CRITERION::INVERSE_GAUSSIAN>> InverseGaussianObjectiveTestF; TEST_P(InverseGaussianObjectiveTestF, InverseGaussianObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, InverseGaussianObjectiveTestF, ::testing::ValuesIn(invgauss_objective_test_parameters)); -typedef ObjectiveTest, +typedef ObjectiveTest, CRITERION::INVERSE_GAUSSIAN>> WeightedInverseGaussianObjectiveTestD; TEST_P(WeightedInverseGaussianObjectiveTestD, InverseGaussianObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedInverseGaussianObjectiveTestD, ::testing::ValuesIn(invgauss_objective_test_parameters)); -typedef ObjectiveTest, - CRITERION::INVERSE_GAUSSIAN>> +typedef ObjectiveTest< + ObjectiveTestConfig, CRITERION::INVERSE_GAUSSIAN>> WeightedInverseGaussianObjectiveTestF; TEST_P(WeightedInverseGaussianObjectiveTestF, InverseGaussianObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2618,28 +2555,28 @@ INSTANTIATE_TEST_CASE_P(RfTests, // entropy objective test typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::ENTROPY>> + ObjectiveTestConfig, CRITERION::ENTROPY>> EntropyObjectiveTestD; TEST_P(EntropyObjectiveTestD, entropyObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, EntropyObjectiveTestD, ::testing::ValuesIn(entropy_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::ENTROPY>> + ObjectiveTestConfig, CRITERION::ENTROPY>> EntropyObjectiveTestF; TEST_P(EntropyObjectiveTestF, entropyObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, EntropyObjectiveTestF, ::testing::ValuesIn(entropy_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::ENTROPY>> + ObjectiveTestConfig, CRITERION::ENTROPY>> WeightedEntropyObjectiveTestD; TEST_P(WeightedEntropyObjectiveTestD, entropyObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedEntropyObjectiveTestD, ::testing::ValuesIn(entropy_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::ENTROPY>> + ObjectiveTestConfig, CRITERION::ENTROPY>> WeightedEntropyObjectiveTestF; TEST_P(WeightedEntropyObjectiveTestF, entropyObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2648,28 +2585,28 @@ INSTANTIATE_TEST_CASE_P(RfTests, // gini objective test typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GINI>> + ObjectiveTestConfig, CRITERION::GINI>> GiniObjectiveTestD; TEST_P(GiniObjectiveTestD, giniObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, GiniObjectiveTestD, ::testing::ValuesIn(gini_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GINI>> + ObjectiveTestConfig, CRITERION::GINI>> GiniObjectiveTestF; TEST_P(GiniObjectiveTestF, giniObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, GiniObjectiveTestF, ::testing::ValuesIn(gini_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GINI>> + ObjectiveTestConfig, CRITERION::GINI>> WeightedGiniObjectiveTestD; TEST_P(WeightedGiniObjectiveTestD, giniObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, WeightedGiniObjectiveTestD, ::testing::ValuesIn(gini_objective_test_parameters)); typedef ObjectiveTest< - ObjectiveTestConfig, CRITERION::GINI>> + ObjectiveTestConfig, CRITERION::GINI>> WeightedGiniObjectiveTestF; TEST_P(WeightedGiniObjectiveTestF, giniObjectiveTest) {} INSTANTIATE_TEST_CASE_P(RfTests, @@ -2704,14 +2641,17 @@ class FeatureSamplingBiasTest : public ::testing::TestWithParamget_stream(); + const auto node_count = ML::narrow_cast(params.n_nodes); + const auto feature_count = ML::narrow_cast(params.n_features); + const auto sampled_col_count = ML::checked_mul(params.n_nodes, params.k); + // Allocate device memory - rmm::device_uvector d_colids(params.n_nodes * params.k, stream); - rmm::device_uvector d_work_items(params.n_nodes, stream); - rmm::device_uvector d_counts(params.n_features, stream); + rmm::device_uvector d_colids(sampled_col_count, stream); + rmm::device_uvector d_work_items(node_count, stream); // Initialize work items on host - std::vector h_work_items(params.n_nodes); - for (int i = 0; i < params.n_nodes; ++i) { + std::vector h_work_items(node_count); + for (std::size_t i = 0; i < node_count; ++i) { h_work_items[i].idx = i; h_work_items[i].depth = 0; h_work_items[i].instances.begin = 0; @@ -2719,58 +2659,35 @@ class FeatureSamplingBiasTest : public ::testing::TestWithParam= 0 && feature_idx < params.n_features) { - EXPECT_FALSE(feature_seen[feature_idx]) << "Node " << node << " has duplicate feature " + const auto feature_pos = ML::narrow_cast(feature_idx); + h_counts[feature_pos]++; + EXPECT_FALSE(feature_seen[feature_pos]) << "Node " << node << " has duplicate feature " << feature_idx << " at positions in sampled set"; - if (!feature_seen[feature_idx]) { - feature_seen[feature_idx] = true; + if (!feature_seen[feature_pos]) { + feature_seen[feature_pos] = true; unique_count++; } } @@ -2795,8 +2714,8 @@ class FeatureSamplingBiasTest : public ::testing::TestWithParam(params.n_nodes, params.k); + double expected_per_feature = double(total_samples) / params.n_features; // Check for feature 0 under-sampling double feature_0_ratio = h_counts[0] / expected_per_feature; @@ -2804,12 +2723,12 @@ class FeatureSamplingBiasTest : public ::testing::TestWithParam