-
-
Notifications
You must be signed in to change notification settings - Fork 8.7k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Split Features into Groups to Compute Histograms in Shared Memory (#5795
- Loading branch information
Showing
6 changed files
with
295 additions
and
42 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,64 @@ | ||
| /*! | ||
| * Copyright 2020 by XGBoost Contributors | ||
| */ | ||
|
|
||
| #include <xgboost/base.h> | ||
| #include <algorithm> | ||
| #include <vector> | ||
|
|
||
| #include "feature_groups.cuh" | ||
|
|
||
| #include "../../common/device_helpers.cuh" | ||
| #include "../../common/hist_util.h" | ||
|
|
||
| namespace xgboost { | ||
| namespace tree { | ||
|
|
||
| FeatureGroups::FeatureGroups(const common::HistogramCuts& cuts, bool is_dense, | ||
| size_t shm_size, size_t bin_size) { | ||
| // Only use a single feature group for sparse matrices. | ||
| bool single_group = !is_dense; | ||
| if (single_group) { | ||
| InitSingle(cuts); | ||
| return; | ||
| } | ||
|
|
||
| std::vector<int>& feature_segments_h = feature_segments.HostVector(); | ||
| std::vector<int>& bin_segments_h = bin_segments.HostVector(); | ||
| feature_segments_h.push_back(0); | ||
| bin_segments_h.push_back(0); | ||
|
|
||
| const std::vector<uint32_t>& cut_ptrs = cuts.Ptrs(); | ||
| int max_shmem_bins = shm_size / bin_size; | ||
| max_group_bins = 0; | ||
|
|
||
| for (size_t i = 2; i < cut_ptrs.size(); ++i) { | ||
| int last_start = bin_segments_h.back(); | ||
| if (cut_ptrs[i] - last_start > max_shmem_bins) { | ||
| feature_segments_h.push_back(i - 1); | ||
| bin_segments_h.push_back(cut_ptrs[i - 1]); | ||
| max_group_bins = std::max(max_group_bins, | ||
| bin_segments_h.back() - last_start); | ||
| } | ||
| } | ||
| feature_segments_h.push_back(cut_ptrs.size() - 1); | ||
| bin_segments_h.push_back(cut_ptrs.back()); | ||
| max_group_bins = std::max(max_group_bins, | ||
| bin_segments_h.back() - | ||
| bin_segments_h[bin_segments_h.size() - 2]); | ||
| } | ||
|
|
||
| void FeatureGroups::InitSingle(const common::HistogramCuts& cuts) { | ||
| std::vector<int>& feature_segments_h = feature_segments.HostVector(); | ||
| feature_segments_h.push_back(0); | ||
| feature_segments_h.push_back(cuts.Ptrs().size() - 1); | ||
|
|
||
| std::vector<int>& bin_segments_h = bin_segments.HostVector(); | ||
| bin_segments_h.push_back(0); | ||
| bin_segments_h.push_back(cuts.TotalBins()); | ||
|
|
||
| max_group_bins = cuts.TotalBins(); | ||
| } | ||
|
|
||
| } // namespace tree | ||
| } // namespace xgboost |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,119 @@ | ||
| /*! | ||
| * Copyright 2020 by XGBoost Contributors | ||
| */ | ||
| #ifndef FEATURE_GROUPS_CUH_ | ||
| #define FEATURE_GROUPS_CUH_ | ||
|
|
||
| #include <xgboost/host_device_vector.h> | ||
| #include <xgboost/span.h> | ||
|
|
||
| namespace xgboost { | ||
|
|
||
| // Forward declarations. | ||
| namespace common { | ||
| class HistogramCuts; | ||
| } // namespace common | ||
|
|
||
| namespace tree { | ||
|
|
||
| /** \brief FeatureGroup is a feature group. It is defined by a range of | ||
| consecutive feature indices, and also contains a range of all bin indices | ||
| associated with those features. */ | ||
| struct FeatureGroup { | ||
| __host__ __device__ FeatureGroup(int start_feature_, int num_features_, | ||
| int start_bin_, int num_bins_) : | ||
| start_feature(start_feature_), num_features(num_features_), | ||
| start_bin(start_bin_), num_bins(num_bins_) {} | ||
| /** The first feature of the group. */ | ||
| int start_feature; | ||
| /** The number of features in the group. */ | ||
| int num_features; | ||
| /** The first bin in the group. */ | ||
| int start_bin; | ||
| /** The number of bins in the group. */ | ||
| int num_bins; | ||
| }; | ||
|
|
||
| /** \brief FeatureGroupsAccessor is a non-owning accessor for FeatureGroups. */ | ||
| struct FeatureGroupsAccessor { | ||
| FeatureGroupsAccessor(common::Span<const int> feature_segments_, | ||
| common::Span<const int> bin_segments_, int max_group_bins_) : | ||
| feature_segments(feature_segments_), bin_segments(bin_segments_), | ||
| max_group_bins(max_group_bins_) {} | ||
|
|
||
| common::Span<const int> feature_segments; | ||
| common::Span<const int> bin_segments; | ||
| int max_group_bins; | ||
|
|
||
| /** \brief Gets the number of feature groups. */ | ||
| __host__ __device__ int NumGroups() const { | ||
| return feature_segments.size() - 1; | ||
| } | ||
|
|
||
| /** \brief Gets the information about a feature group with index i. */ | ||
| __host__ __device__ FeatureGroup operator[](int i) const { | ||
| return {feature_segments[i], feature_segments[i + 1] - feature_segments[i], | ||
| bin_segments[i], bin_segments[i + 1] - bin_segments[i]}; | ||
| } | ||
| }; | ||
|
|
||
| /** \brief FeatureGroups contains information that defines a split of features | ||
| into groups. Bins of a single feature group typically fit into shared | ||
| memory, so the histogram for the features of a single group can be computed | ||
| faster. | ||
| \notes Known limitations: | ||
| - splitting features into groups currently works only for dense matrices, | ||
| where it is easy to get a feature value in a row by its index; for sparse | ||
| matrices, the structure contains only a single group containing all | ||
| features; | ||
| - if a single feature requires more bins than fit into shared memory, the | ||
| histogram is computed in global memory even if there are multiple feature | ||
| groups; note that this is unlikely to occur in practice, as the default | ||
| number of bins per feature is 256, whereas a thread block with 48 KiB | ||
| shared memory can contain 3072 bins if each gradient sum component is a | ||
| 64-bit floating-point value (double) | ||
| */ | ||
| struct FeatureGroups { | ||
| /** Group cuts for features. Size equals to (number of groups + 1). */ | ||
| HostDeviceVector<int> feature_segments; | ||
| /** Group cuts for bins. Size equals to (number of groups + 1) */ | ||
| HostDeviceVector<int> bin_segments; | ||
| /** Maximum number of bins in a group. Useful to compute the amount of dynamic | ||
| shared memory when launching a kernel. */ | ||
| int max_group_bins; | ||
|
|
||
| /** Creates feature groups by splitting features into groups. | ||
| \param cuts Histogram cuts that given the number of bins per feature. | ||
| \param is_dense Whether the data matrix is dense. | ||
| \param shm_size Available size of shared memory per thread block (in | ||
| bytes) used to compute feature groups. | ||
| \param bin_size Size of a single bin of the histogram. */ | ||
| FeatureGroups(const common::HistogramCuts& cuts, bool is_dense, | ||
| size_t shm_size, size_t bin_size); | ||
|
|
||
| /** Creates a single feature group containing all features and bins. | ||
| \notes This is used as a fallback for sparse matrices, and is also useful | ||
| for testing. | ||
| */ | ||
| explicit FeatureGroups(const common::HistogramCuts& cuts) { | ||
| InitSingle(cuts); | ||
| } | ||
|
|
||
| FeatureGroupsAccessor DeviceAccessor(int device) const { | ||
| feature_segments.SetDevice(device); | ||
| bin_segments.SetDevice(device); | ||
| return {feature_segments.ConstDeviceSpan(), bin_segments.ConstDeviceSpan(), | ||
| max_group_bins}; | ||
| } | ||
|
|
||
| private: | ||
| void InitSingle(const common::HistogramCuts& cuts); | ||
| }; | ||
|
|
||
| } // namespace tree | ||
| } // namespace xgboost | ||
|
|
||
| #endif // FEATURE_GROUPS_CUH_ |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.