[EM] Treat ellpack as dense matrix when there's no compression.

src/common/compressed_iterator.h

@@ -1,12 +1,11 @@
 /**
- * Copyright 2017-2023 by XGBoost Contributors
+ * Copyright 2017-2024, XGBoost Contributors
  * \file compressed_iterator.h
  */
 #pragma once
 #include <xgboost/base.h>
 
-#include <algorithm>
-#include <cmath>
+#include <cmath>    // for ceil, log2
 #include <cstddef>  // for size_t
 
 #include "common.h"
@@ -15,9 +14,7 @@
 #include "device_helpers.cuh"
 #endif  // __CUDACC__
 
-namespace xgboost {
-namespace common {
-
+namespace xgboost::common {
 using CompressedByteT = unsigned char;
 
 namespace detail {
@@ -87,13 +84,12 @@ class CompressedBufferWriter {
 
   template <typename T>
   void WriteSymbol(CompressedByteT *buffer, T symbol, size_t offset) {
-    const int bits_per_byte = 8;
+    constexpr std::int32_t kBitsPerByte = 8;
 
     for (size_t i = 0; i < symbol_bits_; i++) {
-      size_t byte_idx = ((offset + 1) * symbol_bits_ - (i + 1)) / bits_per_byte;
+      size_t byte_idx = ((offset + 1) * symbol_bits_ - (i + 1)) / kBitsPerByte;
       byte_idx += detail::kPadding;
-      size_t bit_idx =
-          ((bits_per_byte + i) - ((offset + 1) * symbol_bits_)) % bits_per_byte;
+      size_t bit_idx = ((kBitsPerByte + i) - ((offset + 1) * symbol_bits_)) % kBitsPerByte;
 
       if (detail::CheckBit(symbol, i)) {
         detail::SetBit(&buffer[byte_idx], bit_idx);
@@ -181,16 +177,14 @@ class CompressedIterator {
   typedef value_type reference;             // NOLINT
 
  private:
-  const CompressedByteT *buffer_ {nullptr};
-  size_t symbol_bits_ {0};
+  CompressedByteT const *buffer_{nullptr};
+  bst_idx_t const symbol_bits_{0};
   size_t offset_ {0};
 
  public:
   CompressedIterator() = default;
-  CompressedIterator(const CompressedByteT *buffer, size_t num_symbols)
-      : buffer_(buffer) {
-    symbol_bits_ = detail::SymbolBits(num_symbols);
-  }
+  CompressedIterator(CompressedByteT const *buffer, bst_idx_t num_symbols)
+      : buffer_{buffer}, symbol_bits_{detail::SymbolBits(num_symbols)} {}
 
   XGBOOST_DEVICE reference operator*() const {
     const int bits_per_byte = 8;
@@ -218,5 +212,4 @@ class CompressedIterator {
     return *offset;
   }
 };
-}  // namespace common
-}  // namespace xgboost
+}  // namespace xgboost::common

src/common/hist_util.cuh

@@ -24,29 +24,6 @@
 #include "xgboost/span.h"  // for IterSpan
 
 namespace xgboost::common {
-namespace cuda {
-/**
- * copy and paste of the host version, we can't make it a __host__ __device__ function as
- * the fn might be a host only or device only callable object, which is not allowed by nvcc.
- */
-template <typename Fn>
-auto __device__ DispatchBinType(BinTypeSize type, Fn&& fn) {
-  switch (type) {
-    case kUint8BinsTypeSize: {
-      return fn(uint8_t{});
-    }
-    case kUint16BinsTypeSize: {
-      return fn(uint16_t{});
-    }
-    case kUint32BinsTypeSize: {
-      return fn(uint32_t{});
-    }
-  }
-  SPAN_CHECK(false);
-  return fn(uint32_t{});
-}
-}  // namespace cuda
-
 namespace detail {
 struct EntryCompareOp {
   __device__ bool operator()(const Entry& a, const Entry& b) {

src/common/hist_util.h

@@ -84,6 +84,10 @@ class HistogramCuts {
   [[nodiscard]] bst_bin_t FeatureBins(bst_feature_t feature) const {
     return cut_ptrs_.ConstHostVector().at(feature + 1) - cut_ptrs_.ConstHostVector()[feature];
   }
+  [[nodiscard]] bst_feature_t NumFeatures() const {
+    CHECK_EQ(this->min_vals_.Size(), this->cut_ptrs_.Size() - 1);
+    return this->min_vals_.Size();
+  }
 
   std::vector<uint32_t> const& Ptrs()      const { return cut_ptrs_.ConstHostVector();   }
   std::vector<float>    const& Values()    const { return cut_values_.ConstHostVector(); }
@@ -101,8 +105,10 @@ class HistogramCuts {
     has_categorical_ = has_cat;
     max_cat_ = max_cat;
   }
-
-  [[nodiscard]] bst_bin_t TotalBins() const { return cut_ptrs_.ConstHostVector().back(); }
+  /**
+   * @brief The total number of histogram bins (excluding min values.)
+   */
+  [[nodiscard]] bst_bin_t TotalBins() const { return this->cut_values_.Size(); }
 
   // Return the index of a cut point that is strictly greater than the input
   // value, or the last available index if none exists

src/data/ellpack_page.cc

@@ -53,12 +53,6 @@ bst_idx_t EllpackPage::Size() const {
                 "EllpackPage is required";
   return 0;
 }
-
-[[nodiscard]] bool EllpackPage::IsDense() const {
-  LOG(FATAL) << "Internal Error: XGBoost is not compiled with CUDA but "
-                "EllpackPage is required";
-  return false;
-}
 }  // namespace xgboost
 
 #endif  // XGBOOST_USE_CUDA