Fix CPU hist init for sparse dataset. (dmlc#4625)

* Fix CPU hist init for sparse dataset.

* Implement sparse histogram cut.
* Allow empty features.

* Fix windows build, don't use sparse in distributed environment.

* Comments.

* Smaller threshold.

* Fix windows omp.

* Fix msvc lambda capture.

* Fix MSVC macro.

* Fix MSVC initialization list.

* Fix MSVC initialization list x2.

* Preserve categorical feature behavior.

* Rename matrix to sparse cuts.
* Reuse UseGroup.
* Check for categorical data when adding cut.

Co-Authored-By: Philip Hyunsu Cho <[email protected]>

* Sanity check.

* Fix comments.

* Fix comment.

Author: trivialfis

.clang-tidy

@@ -1,4 +1,4 @@
-Checks: 'modernize-*,-modernize-make-*,-modernize-use-auto,-modernize-raw-string-literal,google-*,-google-default-arguments,-clang-diagnostic-#pragma-messages,readability-identifier-naming'
+Checks: 'modernize-*,-modernize-make-*,-modernize-use-auto,-modernize-raw-string-literal,-modernize-avoid-c-arrays,google-*,-google-default-arguments,-clang-diagnostic-#pragma-messages,readability-identifier-naming'
 CheckOptions:
   - { key: readability-identifier-naming.ClassCase,           value: CamelCase  }
   - { key: readability-identifier-naming.StructCase,          value: CamelCase  }

include/xgboost/data.h

@@ -437,6 +437,7 @@ class DMatrix {
                        bool load_row_split,
                        const std::string& file_format = "auto",
                        const size_t page_size = kPageSize);
+
   /*!
    * \brief create a new DMatrix, by wrapping a row_iterator, and meta info.
    * \param source The source iterator of the data, the create function takes ownership of the source.

src/CMakeLists.txt

@@ -59,7 +59,7 @@ if (USE_CUDA)
 
   # OpenMP is mandatory for cuda version
   find_package(OpenMP REQUIRED)
-  target_compile_options(objxgboost PRIVATE  
+  target_compile_options(objxgboost PRIVATE
     $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=${OpenMP_CXX_FLAGS}>
   )
 

src/c_api/c_api.cc

@@ -119,10 +119,9 @@ class NativeDataIter : public dmlc::Parser<uint32_t> {
   }
 
   bool Next() override {
-    if ((*next_callback_)(
-            data_handle_,
-            XGBoostNativeDataIterSetData,
-            this) != 0) {
+    if ((*next_callback_)(data_handle_,
+                          XGBoostNativeDataIterSetData,
+                          this) != 0) {
       at_first_ = false;
       return true;
     } else {

src/common/column_matrix.h

@@ -75,7 +75,7 @@ class ColumnMatrix {
   // construct column matrix from GHistIndexMatrix
   inline void Init(const GHistIndexMatrix& gmat,
                    double  sparse_threshold) {
-    const int32_t nfeature = static_cast<int32_t>(gmat.cut.row_ptr.size() - 1);
+    const int32_t nfeature = static_cast<int32_t>(gmat.cut.Ptrs().size() - 1);
     const size_t nrow = gmat.row_ptr.size() - 1;
 
     // identify type of each column
@@ -85,7 +85,7 @@ class ColumnMatrix {
 
     uint32_t max_val = std::numeric_limits<uint32_t>::max();
     for (int32_t fid = 0; fid < nfeature; ++fid) {
-      CHECK_LE(gmat.cut.row_ptr[fid + 1] - gmat.cut.row_ptr[fid], max_val);
+      CHECK_LE(gmat.cut.Ptrs()[fid + 1] - gmat.cut.Ptrs()[fid], max_val);
     }
 
     gmat.GetFeatureCounts(&feature_counts_[0]);
@@ -123,7 +123,7 @@ class ColumnMatrix {
     // store least bin id for each feature
     index_base_.resize(nfeature);
     for (int32_t fid = 0; fid < nfeature; ++fid) {
-      index_base_[fid] = gmat.cut.row_ptr[fid];
+      index_base_[fid] = gmat.cut.Ptrs()[fid];
     }
 
     // pre-fill index_ for dense columns
@@ -150,9 +150,9 @@ class ColumnMatrix {
       size_t fid = 0;
       for (size_t i = ibegin; i < iend; ++i) {
         const uint32_t bin_id = gmat.index[i];
-        while (bin_id >= gmat.cut.row_ptr[fid + 1]) {
-          ++fid;
-        }
+        auto iter = std::upper_bound(gmat.cut.Ptrs().cbegin() + fid,
+                                     gmat.cut.Ptrs().cend(), bin_id);
+        fid = std::distance(gmat.cut.Ptrs().cbegin(), iter) - 1;
         if (type_[fid] == kDenseColumn) {
           uint32_t* begin = &index_[boundary_[fid].index_begin];
           begin[rid] = bin_id - index_base_[fid];

src/common/common.h

@@ -72,6 +72,11 @@ inline std::string ToString(const T& data) {
   return os.str();
 }
 
+template <typename T1, typename T2>
+XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b) {
+  return static_cast<T1>(std::ceil(static_cast<double>(a) / b));
+}
+
 /*
  * Range iterator
  */

src/common/config.h

@@ -30,12 +30,13 @@ class ConfigParser {
    * \param path path to configuration file
    */
   explicit ConfigParser(const std::string& path)
-    : line_comment_regex_("^#"),
+      : path_(path),
+      line_comment_regex_("^#"),
       key_regex_(R"rx(^([^#"'=\r\n\t ]+)[\t ]*=)rx"),
       key_regex_escaped_(R"rx(^(["'])([^"'=\r\n]+)\1[\t ]*=)rx"),
       value_regex_(R"rx(^([^#"'=\r\n\t ]+)[\t ]*(?:#.*){0,1}$)rx"),
-      value_regex_escaped_(R"rx(^(["'])([^"'=\r\n]+)\1[\t ]*(?:#.*){0,1}$)rx"),
-      path_(path) {}
+      value_regex_escaped_(R"rx(^(["'])([^"'=\r\n]+)\1[\t ]*(?:#.*){0,1}$)rx")
+  {}
 
   std::string LoadConfigFile(const std::string& path) {
     std::ifstream fin(path, std::ios_base::in | std::ios_base::binary);
@@ -77,8 +78,6 @@ class ConfigParser {
     content = NormalizeConfigEOL(content);
     std::stringstream ss { content };
     std::vector<std::pair<std::string, std::string>> results;
-    char delimiter = '=';
-    char comment = '#';
     std::string line;
     std::string key, value;
     // Loop over every line of the configuration file

src/common/device_helpers.cuh

@@ -1,5 +1,5 @@
 /*!
- * Copyright 2017 XGBoost contributors
+ * Copyright 2017-2019 XGBoost contributors
  */
 #pragma once
 #include <thrust/device_ptr.h>
@@ -183,11 +183,6 @@ __device__ void BlockFill(IterT begin, size_t n, ValueT value) {
  * Kernel launcher
  */
 
-template <typename T1, typename T2>
-T1 DivRoundUp(const T1 a, const T2 b) {
-  return static_cast<T1>(ceil(static_cast<double>(a) / b));
-}
-
 template <typename L>
 __global__ void LaunchNKernel(size_t begin, size_t end, L lambda) {
   for (auto i : GridStrideRange(begin, end)) {
@@ -211,7 +206,7 @@ inline void LaunchN(int device_idx, size_t n, cudaStream_t stream, L lambda) {
   safe_cuda(cudaSetDevice(device_idx));
 
   const int GRID_SIZE =
-      static_cast<int>(DivRoundUp(n, ITEMS_PER_THREAD * BLOCK_THREADS));
+      static_cast<int>(xgboost::common::DivRoundUp(n, ITEMS_PER_THREAD * BLOCK_THREADS));
   LaunchNKernel<<<GRID_SIZE, BLOCK_THREADS, 0, stream>>>(static_cast<size_t>(0),
                                                          n, lambda);
 }
@@ -619,7 +614,7 @@ struct CubMemory {
     if (this->IsAllocated()) {
       XGBDeviceAllocator<uint8_t> allocator;
       allocator.deallocate(thrust::device_ptr<uint8_t>(static_cast<uint8_t *>(d_temp_storage)),
-        temp_storage_bytes);
+                           temp_storage_bytes);
       d_temp_storage = nullptr;
       temp_storage_bytes = 0;
     }
@@ -738,7 +733,7 @@ void SparseTransformLbs(int device_idx, dh::CubMemory *temp_memory,
   const int BLOCK_THREADS = 256;
   const int ITEMS_PER_THREAD = 1;
   const int TILE_SIZE = BLOCK_THREADS * ITEMS_PER_THREAD;
-  auto num_tiles = dh::DivRoundUp(count + num_segments, BLOCK_THREADS);
+  auto num_tiles = xgboost::common::DivRoundUp(count + num_segments, BLOCK_THREADS);
   CHECK(num_tiles < std::numeric_limits<unsigned int>::max());
 
   temp_memory->LazyAllocate(sizeof(CoordinateT) * (num_tiles + 1));
@@ -1158,7 +1153,7 @@ class AllReducer {
   };
 
   /**
-   * \brief Synchronizes the device 
+   * \brief Synchronizes the device
    *
    * \param device_id Identifier for the device.
    */