Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Fix mixed types in GPU sketching. (#7228)

Browse Source
This commit is contained in:
Jiaming Yuan 2021-09-16 00:10:25 +08:00 committed by GitHub
parent 037dd0820d
commit 2942dc68e4
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
2 changed files with 43 additions and 15 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

32
src/common/hist_util.cu
View File

@ -156,22 +156,24 @@ void RemoveDuplicatedCategories(
auto d_old_column_sizes_scan = dh::ToSpan(column_sizes_scan); auto d_old_column_sizes_scan = dh::ToSpan(column_sizes_scan);
dh::caching_device_vector<SketchContainer::OffsetT> new_cuts_size( dh::caching_device_vector<SketchContainer::OffsetT> new_cuts_size(
info.num_col_ + 1); info.num_col_ + 1);
auto d_new_cuts_size = dh::ToSpan(new_cuts_size);
auto d_new_columns_ptr = dh::ToSpan(new_column_scan);
CHECK_EQ(new_column_scan.size(), new_cuts_size.size()); CHECK_EQ(new_column_scan.size(), new_cuts_size.size());
dh::LaunchN(new_column_scan.size(), [=] __device__(size_t idx) { dh::LaunchN(
d_old_column_sizes_scan[idx] = d_new_columns_ptr[idx]; new_column_scan.size(),
if (idx == d_new_columns_ptr.size() - 1) { [=, d_new_cuts_size = dh::ToSpan(new_cuts_size),
return; d_old_column_sizes_scan = dh::ToSpan(column_sizes_scan),
} d_new_columns_ptr = dh::ToSpan(new_column_scan)] __device__(size_t idx) {
if (IsCat(d_feature_types, idx)) { d_old_column_sizes_scan[idx] = d_new_columns_ptr[idx];
// Cut size is the same as number of categories in input. if (idx == d_new_columns_ptr.size() - 1) {
d_new_cuts_size[idx] = return;
d_new_columns_ptr[idx + 1] - d_new_columns_ptr[idx]; }
} else { if (IsCat(d_feature_types, idx)) {
d_new_cuts_size[idx] = d_cuts_ptr[idx] - d_cuts_ptr[idx]; // Cut size is the same as number of categories in input.
} d_new_cuts_size[idx] =
}); d_new_columns_ptr[idx + 1] - d_new_columns_ptr[idx];
} else {
d_new_cuts_size[idx] = d_cuts_ptr[idx + 1] - d_cuts_ptr[idx];
}
});
// Turn size into ptr. // Turn size into ptr.
thrust::exclusive_scan(thrust::device, new_cuts_size.cbegin(), thrust::exclusive_scan(thrust::device, new_cuts_size.cbegin(),
new_cuts_size.cend(), d_cuts_ptr.data()); new_cuts_size.cend(), d_cuts_ptr.data());

26
tests/cpp/common/test_hist_util.cu
View File

@ -165,6 +165,32 @@ TEST(HistUtil, DeviceSketchCategoricalFeatures) {
TestCategoricalSketch(1000, 256, 32, true); TestCategoricalSketch(1000, 256, 32, true);
} }
void TestMixedSketch() {
size_t n_samples = 1000, n_features = 2, n_categories = 3;
std::vector<float> data(n_samples * n_features);
SimpleLCG gen;
SimpleRealUniformDistribution<float> cat_d{0.0f, float(n_categories)};
SimpleRealUniformDistribution<float> num_d{0.0f, 3.0f};
for (size_t i = 0; i < n_samples * n_features; ++i) {
if (i % 2 == 0) {
data[i] = std::floor(cat_d(&gen));
} else {
data[i] = num_d(&gen);
}
}
auto m = GetDMatrixFromData(data, n_samples, n_features);
m->Info().feature_types.HostVector().push_back(FeatureType::kCategorical);
m->Info().feature_types.HostVector().push_back(FeatureType::kNumerical);
auto cuts = DeviceSketch(0, m.get(), 64);
ASSERT_EQ(cuts.Values().size(), 64 + n_categories);
}
TEST(HistUtil, DeviceSketchMixedFeatures) {
TestMixedSketch();
}
TEST(HistUtil, DeviceSketchMultipleColumns) { TEST(HistUtil, DeviceSketchMultipleColumns) {
int bin_sizes[] = {2, 16, 256, 512}; int bin_sizes[] = {2, 16, 256, 512};
int sizes[] = {100, 1000, 1500}; int sizes[] = {100, 1000, 1500};