Sketching from adapters

[RAMitchell]

Overhaul of GPU sketching code to support sketching on external data.

Todo:
Verify peak memory usage

Some performance charts:

[trivialfis]

I'm a little bit concerned about removing weight, we should do some experiments to estimate the impact on various datasets with different characteristics.

[hcho3]

Weighted quantile sketch is one of prominent contribution of XGBoost paper (2016). Is it necessary to remove it?

[RAMitchell]

We cant currently access weights via the adapter constructor - they get added after in different c_api functions. If we can change the c_api to get weights on dmatrix construction we can do it.

[trivialfis]

Can we create an empty DMatrix handle at first?

[RAMitchell]

Current implementation preserves existing behaviour, using weights to compute quantiles when available from the DMatrix. When we implement DMatrix from adapters we can revisit how to access weights in this case.

[trivialfis]

I defined bst_row_t and bst_feature_t and we should use them more often.

[RAMitchell]

Noted. If something is a sum I think it is reasonable to use size_t. For example a sum of bst_row_t can exceed 32 bits.

[trivialfis]

I don't think the maker function can be a cause of fatal error right? It's just a helper function for type deduction

[trivialfis]

@RAMitchell

This is the whole definition of thrust::make_transform_iterator:

template <class AdaptableUnaryFunction, class Iterator>
inline __host__ __device__
transform_iterator<AdaptableUnaryFunction, Iterator>
make_transform_iterator(Iterator it, AdaptableUnaryFunction fun)
{
  return transform_iterator<AdaptableUnaryFunction, Iterator>(it, fun);
} // end make_transform_iterator

Other than it's a host device function, there's no difference with your implementation. Could you take a deeper look into the fatal error you encountered? I think it's worth the debugging effort.

[RAMitchell]

My implementation explicitly specifies the return type. The issue has always beem automatic return type deduction on msvc. I can look further.

[trivialfis]

Yup thanks! You can just print the return type by debugger or use typeid(RetType).name(). If that's the root cause than there's nothing we can do. But I'm still interested in what other type can MSVC deduce. ;-) It's just an unification algorithm, should not be way too weird.

[RAMitchell]

These are the two return types for my version and thrust:

thrust::transform_iterator<__nv_dl_wrapper_t<__nv_dl_tag<void (__cdecl*)(xgboost::data::CupyAdapter *,unsigned __int64,unsigned __int64,float,xgboost::common::SketchContainer *,int),&xgboost::common::ProcessBatch<xgboost::data::CupyAdapter>,1>,xgboost::data::CupyAdapterBatch const >,thrust::counting_iterator<unsigned __int64,thrust::use_default,thrust::use_default,thrust::use_default>,xgboost::data::COOTuple,thrust::use_default>

thrust::transform_iterator<__nv_dl_wrapper_t<__nv_dl_tag<void (__cdecl*)(xgboost::data::CupyAdapter *,unsigned __int64,unsigned __int64,float,xgboost::common::SketchContainer *,int),&xgboost::common::ProcessBatch<xgboost::data::CupyAdapter>,2>,xgboost::data::CupyAdapterBatch const >,thrust::counting_iterator<unsigned __int64,thrust::use_default,thrust::use_default,thrust::use_default>,thrust::use_default,thrust::use_default>

[trivialfis]

Let me take another look today. It seems weird that the GetRowStride is defined twice and have to pull all data to host.