1 #include <gtest/gtest.h>
2 #include <vector>
3 #include "../../helpers.h"
4 #include "../../../../src/common/categorical.h"
5 #include "../../../../src/tree/gpu_hist/row_partitioner.cuh"
6 #include "../../../../src/tree/gpu_hist/histogram.cuh"
7
8 namespace xgboost {
9 namespace tree {
10
11 template <typename Gradient>
TestDeterministicHistogram(bool is_dense,int shm_size)12 void TestDeterministicHistogram(bool is_dense, int shm_size) {
13 size_t constexpr kBins = 256, kCols = 120, kRows = 16384, kRounds = 16;
14 float constexpr kLower = -1e-2, kUpper = 1e2;
15
16 float sparsity = is_dense ? 0.0f : 0.5f;
17 auto matrix = RandomDataGenerator(kRows, kCols, sparsity).GenerateDMatrix();
18 BatchParam batch_param{0, static_cast<int32_t>(kBins)};
19
20 for (auto const& batch : matrix->GetBatches<EllpackPage>(batch_param)) {
21 auto* page = batch.Impl();
22
23 tree::RowPartitioner row_partitioner(0, kRows);
24 auto ridx = row_partitioner.GetRows(0);
25
26 int num_bins = kBins * kCols;
27 dh::device_vector<Gradient> histogram(num_bins);
28 auto d_histogram = dh::ToSpan(histogram);
29 auto gpair = GenerateRandomGradients(kRows, kLower, kUpper);
30 gpair.SetDevice(0);
31
32 FeatureGroups feature_groups(page->Cuts(), page->is_dense, shm_size,
33 sizeof(Gradient));
34
35 auto rounding = CreateRoundingFactor<Gradient>(gpair.DeviceSpan());
36 BuildGradientHistogram(page->GetDeviceAccessor(0),
37 feature_groups.DeviceAccessor(0), gpair.DeviceSpan(),
38 ridx, d_histogram, rounding);
39
40 std::vector<Gradient> histogram_h(num_bins);
41 dh::safe_cuda(cudaMemcpy(histogram_h.data(), d_histogram.data(),
42 num_bins * sizeof(Gradient),
43 cudaMemcpyDeviceToHost));
44
45 for (size_t i = 0; i < kRounds; ++i) {
46 dh::device_vector<Gradient> new_histogram(num_bins);
47 auto d_new_histogram = dh::ToSpan(new_histogram);
48
49 auto rounding = CreateRoundingFactor<Gradient>(gpair.DeviceSpan());
50 BuildGradientHistogram(page->GetDeviceAccessor(0),
51 feature_groups.DeviceAccessor(0),
52 gpair.DeviceSpan(), ridx, d_new_histogram,
53 rounding);
54
55 std::vector<Gradient> new_histogram_h(num_bins);
56 dh::safe_cuda(cudaMemcpy(new_histogram_h.data(), d_new_histogram.data(),
57 num_bins * sizeof(Gradient),
58 cudaMemcpyDeviceToHost));
59 for (size_t j = 0; j < new_histogram_h.size(); ++j) {
60 ASSERT_EQ(new_histogram_h[j].GetGrad(), histogram_h[j].GetGrad());
61 ASSERT_EQ(new_histogram_h[j].GetHess(), histogram_h[j].GetHess());
62 }
63 }
64
65 {
66 auto gpair = GenerateRandomGradients(kRows, kLower, kUpper);
67 gpair.SetDevice(0);
68
69 // Use a single feature group to compute the baseline.
70 FeatureGroups single_group(page->Cuts());
71
72 dh::device_vector<Gradient> baseline(num_bins);
73 BuildGradientHistogram(page->GetDeviceAccessor(0),
74 single_group.DeviceAccessor(0),
75 gpair.DeviceSpan(), ridx, dh::ToSpan(baseline),
76 rounding);
77
78 std::vector<Gradient> baseline_h(num_bins);
79 dh::safe_cuda(cudaMemcpy(baseline_h.data(), baseline.data().get(),
80 num_bins * sizeof(Gradient),
81 cudaMemcpyDeviceToHost));
82
83 for (size_t i = 0; i < baseline.size(); ++i) {
84 EXPECT_NEAR(baseline_h[i].GetGrad(), histogram_h[i].GetGrad(),
85 baseline_h[i].GetGrad() * 1e-3);
86 }
87 }
88 }
89 }
90
TEST(Histogram,GPUDeterministic)91 TEST(Histogram, GPUDeterministic) {
92 std::vector<bool> is_dense_array{false, true};
93 std::vector<int> shm_sizes{48 * 1024, 64 * 1024, 160 * 1024};
94 for (bool is_dense : is_dense_array) {
95 for (int shm_size : shm_sizes) {
96 TestDeterministicHistogram<GradientPair>(is_dense, shm_size);
97 TestDeterministicHistogram<GradientPairPrecise>(is_dense, shm_size);
98 }
99 }
100 }
101
OneHotEncodeFeature(std::vector<float> x,size_t num_cat)102 std::vector<float> OneHotEncodeFeature(std::vector<float> x, size_t num_cat) {
103 std::vector<float> ret(x.size() * num_cat, 0);
104 size_t n_rows = x.size();
105 for (size_t r = 0; r < n_rows; ++r) {
106 bst_cat_t cat = common::AsCat(x[r]);
107 ret.at(num_cat * r + cat) = 1;
108 }
109 return ret;
110 }
111
112 // Test 1 vs rest categorical histogram is equivalent to one hot encoded data.
TestGPUHistogramCategorical(size_t num_categories)113 void TestGPUHistogramCategorical(size_t num_categories) {
114 size_t constexpr kRows = 340;
115 size_t constexpr kBins = 256;
116 auto x = GenerateRandomCategoricalSingleColumn(kRows, num_categories);
117 auto cat_m = GetDMatrixFromData(x, kRows, 1);
118 cat_m->Info().feature_types.HostVector().push_back(FeatureType::kCategorical);
119 BatchParam batch_param{0, static_cast<int32_t>(kBins)};
120 tree::RowPartitioner row_partitioner(0, kRows);
121 auto ridx = row_partitioner.GetRows(0);
122 dh::device_vector<GradientPairPrecise> cat_hist(num_categories);
123 auto gpair = GenerateRandomGradients(kRows, 0, 2);
124 gpair.SetDevice(0);
125 auto rounding = CreateRoundingFactor<GradientPairPrecise>(gpair.DeviceSpan());
126 // Generate hist with cat data.
127 for (auto const &batch : cat_m->GetBatches<EllpackPage>(batch_param)) {
128 auto* page = batch.Impl();
129 FeatureGroups single_group(page->Cuts());
130 BuildGradientHistogram(page->GetDeviceAccessor(0),
131 single_group.DeviceAccessor(0),
132 gpair.DeviceSpan(), ridx, dh::ToSpan(cat_hist),
133 rounding);
134 }
135
136 // Generate hist with one hot encoded data.
137 auto x_encoded = OneHotEncodeFeature(x, num_categories);
138 auto encode_m = GetDMatrixFromData(x_encoded, kRows, num_categories);
139 dh::device_vector<GradientPairPrecise> encode_hist(2 * num_categories);
140 for (auto const &batch : encode_m->GetBatches<EllpackPage>(batch_param)) {
141 auto* page = batch.Impl();
142 FeatureGroups single_group(page->Cuts());
143 BuildGradientHistogram(page->GetDeviceAccessor(0),
144 single_group.DeviceAccessor(0),
145 gpair.DeviceSpan(), ridx, dh::ToSpan(encode_hist),
146 rounding);
147 }
148
149 std::vector<GradientPairPrecise> h_cat_hist(cat_hist.size());
150 thrust::copy(cat_hist.begin(), cat_hist.end(), h_cat_hist.begin());
151 auto cat_sum = std::accumulate(h_cat_hist.begin(), h_cat_hist.end(), GradientPairPrecise{});
152
153 std::vector<GradientPairPrecise> h_encode_hist(encode_hist.size());
154 thrust::copy(encode_hist.begin(), encode_hist.end(), h_encode_hist.begin());
155
156 for (size_t c = 0; c < num_categories; ++c) {
157 auto zero = h_encode_hist[c * 2];
158 auto one = h_encode_hist[c * 2 + 1];
159
160 auto chosen = h_cat_hist[c];
161 auto not_chosen = cat_sum - chosen;
162
163 ASSERT_LE(RelError(zero.GetGrad(), not_chosen.GetGrad()), kRtEps);
164 ASSERT_LE(RelError(zero.GetHess(), not_chosen.GetHess()), kRtEps);
165
166 ASSERT_LE(RelError(one.GetGrad(), chosen.GetGrad()), kRtEps);
167 ASSERT_LE(RelError(one.GetHess(), chosen.GetHess()), kRtEps);
168 }
169 }
170
TEST(Histogram,GPUHistCategorical)171 TEST(Histogram, GPUHistCategorical) {
172 for (size_t num_categories = 2; num_categories < 8; ++num_categories) {
173 TestGPUHistogramCategorical(num_categories);
174 }
175 }
176 } // namespace tree
177 } // namespace xgboost
178