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