Revert #5755

[hcho3]

Temporarily revert #5755 to prevent OOM error in the e2e mortgage notebook.

[trivialfis]

I'm still looking into it. Expect some results today. Reverting it might break quantile dmatrix completely.

[hcho3]

@trivialfis At least it does not break the gtest. Let's see whether the Python tests would pass too.

[RAMitchell]

The relevant PR removes the specialised copy for column wise data and instead uses this:

// Here the data is already correctly ordered and simply needs to be compacted
// to remove missing data
template <typename AdapterT>
void CopyDataToDMatrix(AdapterT* adapter, common::Span<Entry> data,
                       int device_idx, float missing,
                       common::Span<size_t> row_ptr) {
  auto& batch = adapter->Value();
  auto transform_f = [=] __device__(size_t idx) {
    const auto& e = batch.GetElement(idx);
    return Entry(e.column_idx, e.value);
  };  // NOLINT
  auto counting = thrust::make_counting_iterator(0llu);
  thrust::transform_iterator<decltype(transform_f), decltype(counting), Entry>
      transform_iter(counting, transform_f);
  dh::XGBCachingDeviceAllocator<char> alloc;
  thrust::copy_if(
      thrust::cuda::par(alloc), transform_iter, transform_iter + batch.Size(),
      thrust::device_pointer_cast(data.data()), IsValidFunctor(missing));
}

I wonder if the copy_if is responsible for the extra memory allocation.

[RAMitchell]

Thrust is using n words to store predicates for the output location of copy_if: https://github.com/thrust/thrust/blob/94e0b1c63fadd5c574c45f05267e45ac136e78ef/thrust/system/detail/generic/copy_if.inl#L64

We can write this more efficiently by redirecting the output of an exclusive scan purely with iterators.

Does someone want to try this? Otherwise I can implement myself.

[trivialfis]

Let me give a try.

[codecov-commenter]

Codecov Report

Merging #6145 into master will not change coverage.
The diff coverage is n/a.

@@           Coverage Diff           @@
##           master    #6145   +/-   ##
=======================================
  Coverage   78.83%   78.83%           
=======================================
  Files          12       12           
  Lines        3090     3090           
=======================================
  Hits         2436     2436           
  Misses        654      654           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 33d80ff...52f4cce. Read the comment docs.

[trivialfis]

I don't think this is necessary anymore. @hcho3 Thanks for bisecting the code.