From 2fb7c8a20139b1ed21f01ff7fc08474c2db54e58 Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 26 Oct 2021 21:13:34 +0800 Subject: [PATCH] Implement typed and type erased tensor. * Use typed tensor for storing meta info like base margin. * Extend the array interface handler to multi-dim. Implement a general array view. * Replace existing matrix and vector view. lint. Remove const too. Doc/Test. Include. Use it in AUC. Win build. Use int32_t. Use integral. force the same type. Use constexpr for old nvcc. Test for empty tensor. Rename to view. Format. Better document and perf. Address reviewer's comment. tidy. Implement a general array view. * Replace existing matrix and vector view. lint. Remove const too. Doc/Test. Include. Use it in AUC. Win build. Use int32_t. Use integral. force the same type. Use constexpr for old nvcc. Test for empty tensor. Rename to view. Format. Prototype. Move string view. Compile on CPU. Some fixes for GPU compilation. Array interface. Use it in file iter. Cleanup. itemsize. Documents. Cache the pointer. port. cuda compilation. Start working on ari. Add clang-format config. (#7383) Generated using `clang-format -style=google -dump-config > .clang-format`, with column width changed from 80 to 100 to be consistent with existing cpplint check. Define shape and stride. Convert some uses. Prototype for copy tensor info. proto unravel Indexer. Unravel. Cleanup. Fixes. fixe.s WAR. as column vector. Convert some code. some more. some more. Compile. Ensure column vector from the beginning. IO. Add code comments. Test for trivial dimension. Start CPU implementation. Refactor. Dispatch. Compile. --- include/xgboost/c_api.h | 31 +- include/xgboost/data.h | 16 +- include/xgboost/json.h | 46 +-- include/xgboost/linalg.h | 264 ++++++++++++- include/xgboost/string_view.h | 56 +++ python-package/setup.cfg | 5 +- src/c_api/c_api.cc | 3 +- src/common/host_device_vector.cc | 1 + src/common/host_device_vector.cu | 1 + src/data/adapter.h | 82 ++-- src/data/array_interface.h | 351 +++++++++--------- src/data/data.cc | 142 ++++++- src/data/data.cu | 185 +++++---- src/data/device_adapter.cuh | 48 +-- src/data/file_iterator.h | 15 +- src/data/simple_dmatrix.cc | 7 +- src/gbm/gblinear.cc | 23 +- src/predictor/cpu_predictor.cc | 18 +- src/predictor/gpu_predictor.cu | 20 +- tests/cpp/common/test_linalg.cc | 64 +++- tests/cpp/data/test_adapter.cc | 7 +- tests/cpp/data/test_array_interface.cc | 48 ++- tests/cpp/data/test_array_interface.cu | 2 +- .../cpp/data/test_iterative_device_dmatrix.cu | 4 +- tests/cpp/data/test_metainfo.cc | 7 +- tests/cpp/data/test_metainfo.cu | 4 +- tests/cpp/data/test_simple_dmatrix.cc | 10 +- tests/cpp/predictor/test_gpu_predictor.cu | 4 +- 28 files changed, 980 insertions(+), 484 deletions(-) create mode 100644 include/xgboost/string_view.h diff --git a/include/xgboost/c_api.h b/include/xgboost/c_api.h index e935bae4741b..818b2e065927 100644 --- a/include/xgboost/c_api.h +++ b/include/xgboost/c_api.h @@ -240,7 +240,7 @@ XGB_DLL int XGDMatrixCreateFromCudaArrayInterface(char const *data, char const* json_config, DMatrixHandle *out); -/* +/** * ========================== Begin data callback APIs ========================= * * Short notes for data callback @@ -249,9 +249,9 @@ XGB_DLL int XGDMatrixCreateFromCudaArrayInterface(char const *data, * used by JVM packages. It uses `XGBoostBatchCSR` to accept batches for CSR formated * input, and concatenate them into 1 final big CSR. The related functions are: * - * - XGBCallbackSetData - * - XGBCallbackDataIterNext - * - XGDMatrixCreateFromDataIter + * - \ref XGBCallbackSetData + * - \ref XGBCallbackDataIterNext + * - \ref XGDMatrixCreateFromDataIter * * Another set is used by external data iterator. It accept foreign data iterators as * callbacks. There are 2 different senarios where users might want to pass in callbacks @@ -267,17 +267,17 @@ XGB_DLL int XGDMatrixCreateFromCudaArrayInterface(char const *data, * Related functions are: * * # Factory functions - * - `XGDMatrixCreateFromCallback` for external memory - * - `XGDeviceQuantileDMatrixCreateFromCallback` for quantile DMatrix + * - \ref XGDMatrixCreateFromCallback for external memory + * - \ref XGDeviceQuantileDMatrixCreateFromCallback for quantile DMatrix * * # Proxy that callers can use to pass data to XGBoost - * - XGProxyDMatrixCreate - * - XGDMatrixCallbackNext - * - DataIterResetCallback - * - XGProxyDMatrixSetDataCudaArrayInterface - * - XGProxyDMatrixSetDataCudaColumnar - * - XGProxyDMatrixSetDataDense - * - XGProxyDMatrixSetDataCSR + * - \ref XGProxyDMatrixCreate + * - \ref XGDMatrixCallbackNext + * - \ref DataIterResetCallback + * - \ref XGProxyDMatrixSetDataCudaArrayInterface + * - \ref XGProxyDMatrixSetDataCudaColumnar + * - \ref XGProxyDMatrixSetDataDense + * - \ref XGProxyDMatrixSetDataCSR * - ... (data setters) */ @@ -402,7 +402,7 @@ XGB_EXTERN_C typedef void DataIterResetCallback(DataIterHandle handle); // NOLIN * - cache_prefix: The path of cache file, caller must initialize all the directories in this path. * - nthread (optional): Number of threads used for initializing DMatrix. * - * \param out The created external memory DMatrix + * \param[out] out The created external memory DMatrix * * \return 0 when success, -1 when failure happens */ @@ -596,7 +596,8 @@ XGB_DLL int XGDMatrixSetUIntInfo(DMatrixHandle handle, * char const* feat_names [] {"feat_0", "feat_1"}; * XGDMatrixSetStrFeatureInfo(handle, "feature_name", feat_names, 2); * - * // i for integer, q for quantitive. Similarly "int" and "float" are also recognized. + * // i for integer, q for quantitive, c for categorical. Similarly "int" and "float" + * // are also recognized. * char const* feat_types [] {"i", "q"}; * XGDMatrixSetStrFeatureInfo(handle, "feature_type", feat_types, 2); * diff --git a/include/xgboost/data.h b/include/xgboost/data.h index cd000e371332..c3e660580bbe 100644 --- a/include/xgboost/data.h +++ b/include/xgboost/data.h @@ -11,12 +11,14 @@ #include #include #include -#include #include +#include +#include +#include +#include #include #include -#include #include #include #include @@ -67,7 +69,7 @@ class MetaInfo { * if specified, xgboost will start from this init margin * can be used to specify initial prediction to boost from. */ - HostDeviceVector base_margin_; // NOLINT + linalg::Tensor base_margin_; // NOLINT /*! * \brief lower bound of the label, to be used for survival analysis (censored regression) */ @@ -157,7 +159,8 @@ class MetaInfo { * * Right now only 1 column is permitted. */ - void SetInfo(const char* key, std::string const& interface_str); + + void SetInfo(StringView key, StringView interface_str); void GetInfo(char const* key, bst_ulong* out_len, DataType dtype, const void** out_dptr) const; @@ -179,6 +182,9 @@ class MetaInfo { void Extend(MetaInfo const& that, bool accumulate_rows, bool check_column); private: + void SetInfoFromHost(StringView key, StringView interface_str); + void SetInfoFromCUDA(const char* key, std::string const& interface_str); + /*! \brief argsort of labels */ mutable std::vector label_order_cache_; }; @@ -477,7 +483,7 @@ class DMatrix { this->Info().SetInfo(key, dptr, dtype, num); } virtual void SetInfo(const char* key, std::string const& interface_str) { - this->Info().SetInfo(key, interface_str); + this->Info().SetInfo(key, StringView{interface_str}); } /*! \brief meta information of the dataset */ virtual const MetaInfo& Info() const = 0; diff --git a/include/xgboost/json.h b/include/xgboost/json.h index ab6ba6ee1018..c00fc28a2835 100644 --- a/include/xgboost/json.h +++ b/include/xgboost/json.h @@ -4,16 +4,17 @@ #ifndef XGBOOST_JSON_H_ #define XGBOOST_JSON_H_ +#include #include #include -#include +#include +#include #include #include -#include -#include -#include #include +#include +#include namespace xgboost { @@ -298,43 +299,6 @@ class JsonBoolean : public Value { } }; -struct StringView { - private: - using CharT = char; // unsigned char - using Traits = std::char_traits; - CharT const* str_; - size_t size_; - - public: - StringView() = default; - StringView(CharT const* str, size_t size) : str_{str}, size_{size} {} - explicit StringView(std::string const& str): str_{str.c_str()}, size_{str.size()} {} - explicit StringView(CharT const* str) : str_{str}, size_{Traits::length(str)} {} - - CharT const& operator[](size_t p) const { return str_[p]; } - CharT const& at(size_t p) const { // NOLINT - CHECK_LT(p, size_); - return str_[p]; - } - size_t size() const { return size_; } // NOLINT - // Copies a portion of string. Since we don't have std::from_chars and friends here, so - // copying substring is necessary for appending `\0`. It's not too bad since string by - // default has small vector optimization, which is enabled by most if not all modern - // compilers for numeric values. - std::string substr(size_t beg, size_t n) const { // NOLINT - CHECK_LE(beg, size_); - return std::string {str_ + beg, n < (size_ - beg) ? n : (size_ - beg)}; - } - CharT const* c_str() const { return str_; } // NOLINT - - CharT const* cbegin() const { return str_; } // NOLINT - CharT const* cend() const { return str_ + size(); } // NOLINT - CharT const* begin() const { return str_; } // NOLINT - CharT const* end() const { return str_ + size(); } // NOLINT -}; - -std::ostream &operator<<(std::ostream &os, StringView const v); - /*! * \brief Data structure representing JSON format. * diff --git a/include/xgboost/linalg.h b/include/xgboost/linalg.h index a801228e903e..d356ab690465 100644 --- a/include/xgboost/linalg.h +++ b/include/xgboost/linalg.h @@ -6,8 +6,10 @@ #ifndef XGBOOST_LINALG_H_ #define XGBOOST_LINALG_H_ +#include #include -#include +#include +#include #include #include @@ -19,6 +21,17 @@ namespace xgboost { namespace linalg { namespace detail { + +struct ArrayInterfaceHandler { + // FIXME: Duplicated function. + template + static constexpr char TypeChar() { + return (std::is_floating_point::value + ? 'f' + : (std::is_integral::value ? (std::is_signed::value ? 'i' : 'u') : '\0')); + } +}; + template constexpr size_t Offset(S (&strides)[D], size_t n, size_t dim, Head head) { assert(dim < D); @@ -31,6 +44,14 @@ constexpr size_t Offset(S (&strides)[D], size_t n, size_t dim, Head head, Tail & return Offset(strides, n + (head * strides[dim]), dim + 1, rest...); } +template +constexpr void CalcStride(size_t (&shape)[D], size_t (&stride)[D]) { + stride[D - 1] = 1; + for (int32_t s = D - 2; s >= 0; --s) { + stride[s] = shape[s + 1] * stride[s + 1]; + } +} + struct AllTag {}; struct IntTag {}; @@ -71,6 +92,69 @@ XGBOOST_DEVICE constexpr auto UnrollLoop(Fn fn) { fn(i); } } + +inline XGBOOST_DEVICE int Popc(uint32_t v) { +#if defined(__CUDA_ARCH__) + return __popc(v); +#else + return __builtin_popcountl(v); +#endif +} + +inline XGBOOST_DEVICE int Popc(uint64_t v) { +#if defined(__CUDA_ARCH__) + return __popcll(v); +#else + return __builtin_popcountll(v); +#endif +} + +template +constexpr auto Arr2Tup(T (&arr)[N], std::index_sequence) { + return std::make_tuple(arr[Idx]...); +} + +template +constexpr auto Arr2Tup(T (&arr)[N]) { + return Arr2Tup(arr, std::make_index_sequence{}); +} + +template +XGBOOST_DEVICE auto Unravel(size_t idx, common::Span shape) { + size_t index[D]{0}; + auto impl = [&](auto i) { + auto a = i; + for (int dim = D; --dim > 0;) { + auto s = static_cast(shape[dim]); + if (s & (s - 1)) { + auto t = a / s; + index[dim] = a - t * s; + a = t; + } else { // exp of 2 + index[dim] = a & (s - 1); + a >>= Popc(s - 1); + } + } + index[0] = a; + }; + if (idx > std::numeric_limits::max()) { + impl(static_cast(idx)); + } else { + impl(static_cast(idx)); + } + return Arr2Tup(index); +} + +template +XGBOOST_DEVICE auto constexpr Apply(Fn&& f, Tup&& t, std::index_sequence) { + return f(std::get(t)...); +} + +template +XGBOOST_DEVICE auto constexpr Apply(Fn&& f, Tup&& t) { + constexpr auto kSize = std::tuple_size::value; + return Apply(std::forward(f), std::forward(t), std::make_index_sequence{}); +} } // namespace detail /** @@ -96,7 +180,7 @@ class TensorView { StrideT stride_{1}; ShapeT shape_{0}; common::Span data_; - T* ptr_{nullptr}; // pointer of data_ to avoid bound check. + T *ptr_{nullptr}; // pointer of data_ to avoid bound check. size_t size_{0}; int32_t device_{-1}; @@ -174,10 +258,8 @@ class TensorView { shape_[i] = 1; } // stride - stride_[kDim - 1] = 1; - for (int32_t s = kDim - 2; s >= 0; --s) { - stride_[s] = shape_[s + 1] * stride_[s + 1]; - } + detail::CalcStride(shape_, stride_); + // size this->CalcSize(); }; @@ -277,8 +359,52 @@ class TensorView { XGBOOST_DEVICE auto end() { return data_.end(); } // NOLINT XGBOOST_DEVICE size_t Size() const { return size_; } + /** + * \brief Obtain the raw data. + */ XGBOOST_DEVICE auto Values() const { return data_; } + /** + * \brief Obtain the CUDA device ordinal. + */ XGBOOST_DEVICE auto DeviceIdx() const { return device_; } + + /** + * \brief Array Interface defined by + * numpy. + * + * `stream` is optionally included when data is on CUDA device. + */ + std::string ArrayInterface() const { + Json array_interface{Object{}}; + array_interface["data"] = std::vector(2); + array_interface["data"][0] = Integer(reinterpret_cast(data_.data())); + array_interface["data"][1] = Boolean{false}; + if (this->DeviceIdx() >= 0) { + // Change this once we have different CUDA stream. + array_interface["stream"] = Integer(0); + } + std::vector shape(Shape().size()); + std::vector stride(Stride().size()); + for (size_t i = 0; i < Shape().size(); ++i) { + shape[i] = Integer(Shape(i)); + stride[i] = Integer(Stride(i) * sizeof(T)); + } + array_interface["shape"] = Array{shape}; + array_interface["stride"] = Array{stride}; + array_interface["version"] = 3; + + char constexpr kT = detail::ArrayInterfaceHandler::TypeChar(); + static_assert(kT != '\0', ""); + if (DMLC_LITTLE_ENDIAN) { + array_interface["typestr"] = String{"<" + (kT + std::to_string(sizeof(T)))}; + } else { + array_interface["typestr"] = String{">" + (kT + std::to_string(sizeof(T)))}; + } + + std::string str; + Json::Dump(array_interface, &str); + return str; + } }; /** @@ -289,6 +415,19 @@ class TensorView { template using VectorView = TensorView; +/** + * \brief Create a vector view from contigious memory. + * + * \param ptr Pointer to the contigious memory. + * \param s Size of the vector. + * \param device (optional) Device ordinal, default to be host. + */ +template +auto MakeVec(T *ptr, size_t s, int32_t device = -1) { + using U = std::remove_const_t> const; + return linalg::TensorView{{ptr, s}, {s}, device}; +} + /** * \brief A view over a matrix, specialization of Tensor. * @@ -296,6 +435,119 @@ using VectorView = TensorView; */ template using MatrixView = TensorView; + +/** + * \brief A tensor storage. To use it for other functionality like slicing one needs to + * obtain a view first. This way we can use it on both host and device. + */ +template +class Tensor { + public: + using ShapeT = size_t[kDim]; + using StrideT = ShapeT; + + private: + HostDeviceVector data_; + ShapeT shape_{0}; + + public: + Tensor() = default; + + /** + * \brief Create a tensor with shape and device ordinal. The storage is initialized + * automatically. + * + * See \ref TensorView for parameters of this constructor. + */ + template + explicit Tensor(I const (&shape)[D], int32_t device) { + // No device unroll as this is a host only function. + std::copy(shape, shape + D, shape_); + for (auto i = D; i < kDim; ++i) { + shape_[i] = 1; + } + auto size = detail::CalcSize(shape_); + if (device >= 0) { + data_.SetDevice(device); + } + data_.Resize(size); + if (device >= 0) { + data_.DevicePointer(); // Pull to device + } + } + /** + * Initialize from 2 host iterators. + */ + template + explicit Tensor(It begin, It end, I const (&shape)[D], int32_t device) { + // shape + detail::UnrollLoop([&](auto i) { shape_[i] = shape[i]; }); + for (auto i = D; i < kDim; ++i) { + shape_[i] = 1; + } + auto &h_vec = data_.HostVector(); + h_vec.insert(h_vec.begin(), begin, end); + if (device >= 0) { + data_.SetDevice(device); + data_.DevicePointer(); // Pull to device; + } + } + /** + * \brief Get a \ref TensorView for this tensor. + */ + TensorView View(int32_t device) { + if (device >= 0) { + data_.SetDevice(device); + auto span = data_.DeviceSpan(); + return {span, shape_, device}; + } else { + auto span = data_.HostSpan(); + return {span, shape_, device}; + } + } + TensorView View(int32_t device) const { + if (device >= 0) { + data_.SetDevice(device); + auto span = data_.ConstDeviceSpan(); + return {span, shape_, device}; + } else { + auto span = data_.ConstHostSpan(); + return {span, shape_, device}; + } + } + + size_t Size() const { return data_.Size(); } + auto Shape() const { return common::Span{shape_}; } + auto Shape(size_t i) const { return shape_[i]; } + + HostDeviceVector *Data() { return &data_; } + HostDeviceVector const *Data() const { return &data_; } + + /** + * \brief Visitor function for modification that changes shape and data. + * + * \tparam Fn function that takes a pointer to `HostDeviceVector` and a static sized + * span as parameters. + */ + template + void ModifyInplace(Fn &&fn) { + fn(this->Data(), common::Span{this->shape_}); + CHECK_EQ(this->Data()->Size(), detail::CalcSize(this->shape_)) + << "Inconsistent size after modification."; + } +}; + +// Only first axis is support for now. +template +void Stack(Tensor *l, Tensor const &r) { + if (r.Data()->DeviceIdx() >= 0) { + l->Data()->SetDevice(r.Data()->DeviceIdx()); + } + l->ModifyInplace([&](HostDeviceVector *data, common::Span shape) { + data->Extend(*r.Data()); + shape[0] = l->Shape(0) + r.Shape(0); + }); +} } // namespace linalg } // namespace xgboost #endif // XGBOOST_LINALG_H_ diff --git a/include/xgboost/string_view.h b/include/xgboost/string_view.h new file mode 100644 index 000000000000..a548b59f7a05 --- /dev/null +++ b/include/xgboost/string_view.h @@ -0,0 +1,56 @@ +#ifndef XGBOOST_STRING_VIEW_H +#define XGBOOST_STRING_VIEW_H +#include +#include + +namespace xgboost { +struct StringView { + private: + using CharT = char; // unsigned char + using Traits = std::char_traits; + CharT const* str_{nullptr}; + size_t size_{0}; + + public: + constexpr StringView() = default; + constexpr StringView(CharT const* str, size_t size) : str_{str}, size_{size} {} + explicit StringView(std::string const& str): str_{str.c_str()}, size_{str.size()} {} + StringView(CharT const* str) : str_{str}, size_{Traits::length(str)} {} // NOLINT + + CharT const& operator[](size_t p) const { return str_[p]; } + CharT const& at(size_t p) const { // NOLINT + CHECK_LT(p, size_); + return str_[p]; + } + constexpr size_t size() const { return size_; } // NOLINT + // Copies a portion of string. Since we don't have std::from_chars and friends here, so + // copying substring is necessary for appending `\0`. It's not too bad since string by + // default has small vector optimization, which is enabled by most if not all modern + // compilers for numeric values. + std::string substr(size_t beg, size_t n) const { // NOLINT + CHECK_LE(beg, size_); + return std::string {str_ + beg, n < (size_ - beg) ? n : (size_ - beg)}; + } + CharT const* c_str() const { return str_; } // NOLINT + + constexpr CharT const* cbegin() const { return str_; } // NOLINT + constexpr CharT const* cend() const { return str_ + size(); } // NOLINT + constexpr CharT const* begin() const { return str_; } // NOLINT + constexpr CharT const* end() const { return str_ + size(); } // NOLINT +}; +std::ostream &operator<<(std::ostream &os, StringView const v); + +constexpr bool operator==(StringView l, StringView r) { + if (l.size() != r.size()) { + return false; + } + for (auto l_beg = l.cbegin(), r_beg = r.cbegin(); l_beg != l.cend(); + ++l_beg, ++r_beg) { + if (*l_beg != *r_beg) { + return false; + } + } + return true; +} +} // namespace xgboost +#endif // XGBOOST_STRING_VIEW_H diff --git a/python-package/setup.cfg b/python-package/setup.cfg index 9337cad1d6ec..51a664390411 100644 --- a/python-package/setup.cfg +++ b/python-package/setup.cfg @@ -4,4 +4,7 @@ description-file = README.rst [mypy] ignore_missing_imports = True disallow_untyped_defs = True -follow_imports = silent \ No newline at end of file +follow_imports = silent + +[flake8] +max-line-length=90 \ No newline at end of file diff --git a/src/c_api/c_api.cc b/src/c_api/c_api.cc index 42b9cad6fb4a..816c62e67237 100644 --- a/src/c_api/c_api.cc +++ b/src/c_api/c_api.cc @@ -426,8 +426,7 @@ XGB_DLL int XGDMatrixSetInfoFromInterface(DMatrixHandle handle, char const* interface_c_str) { API_BEGIN(); CHECK_HANDLE(); - static_cast*>(handle) - ->get()->Info().SetInfo(field, interface_c_str); + static_cast *>(handle)->get()->Info().SetInfo(field, interface_c_str); API_END(); } diff --git a/src/common/host_device_vector.cc b/src/common/host_device_vector.cc index cb02b22f0ca3..3a4a59db7881 100644 --- a/src/common/host_device_vector.cc +++ b/src/common/host_device_vector.cc @@ -170,6 +170,7 @@ void HostDeviceVector::SetDevice(int) const {} // explicit instantiations are required, as HostDeviceVector isn't header-only template class HostDeviceVector; +template class HostDeviceVector; template class HostDeviceVector; template class HostDeviceVector; // bst_node_t template class HostDeviceVector; diff --git a/src/common/host_device_vector.cu b/src/common/host_device_vector.cu index 8287cb24a1bd..456c60a67071 100644 --- a/src/common/host_device_vector.cu +++ b/src/common/host_device_vector.cu @@ -398,6 +398,7 @@ void HostDeviceVector::Resize(size_t new_size, T v) { // explicit instantiations are required, as HostDeviceVector isn't header-only template class HostDeviceVector; +template class HostDeviceVector; template class HostDeviceVector; template class HostDeviceVector; // bst_node_t template class HostDeviceVector; diff --git a/src/data/adapter.h b/src/data/adapter.h index 27da8c6e3b36..9f9e2580ee32 100644 --- a/src/data/adapter.h +++ b/src/data/adapter.h @@ -254,20 +254,20 @@ class ArrayAdapterBatch : public detail::NoMetaInfo { static constexpr bool kIsRowMajor = true; private: - ArrayInterface array_interface_; + ArrayInterface<2> array_interface_; class Line { - ArrayInterface array_interface_; + ArrayInterface<2> array_interface_; size_t ridx_; public: - Line(ArrayInterface array_interface, size_t ridx) + Line(ArrayInterface<2> array_interface, size_t ridx) : array_interface_{std::move(array_interface)}, ridx_{ridx} {} - size_t Size() const { return array_interface_.num_cols; } + size_t Size() const { return array_interface_.Shape(1); } COOTuple GetElement(size_t idx) const { - return {ridx_, idx, array_interface_.GetElement(ridx_, idx)}; + return {ridx_, idx, array_interface_(ridx_, idx)}; } }; @@ -277,11 +277,11 @@ class ArrayAdapterBatch : public detail::NoMetaInfo { return Line{array_interface_, idx}; } - size_t NumRows() const { return array_interface_.num_rows; } - size_t NumCols() const { return array_interface_.num_cols; } + size_t NumRows() const { return array_interface_.Shape(0); } + size_t NumCols() const { return array_interface_.Shape(1); } size_t Size() const { return this->NumRows(); } - explicit ArrayAdapterBatch(ArrayInterface array_interface) + explicit ArrayAdapterBatch(ArrayInterface<2> array_interface) : array_interface_{std::move(array_interface)} {} }; @@ -294,43 +294,43 @@ class ArrayAdapter : public detail::SingleBatchDataIter { public: explicit ArrayAdapter(StringView array_interface) { auto j = Json::Load(array_interface); - array_interface_ = ArrayInterface(get(j)); + array_interface_ = ArrayInterface<2>(get(j)); batch_ = ArrayAdapterBatch{array_interface_}; } ArrayAdapterBatch const& Value() const override { return batch_; } - size_t NumRows() const { return array_interface_.num_rows; } - size_t NumColumns() const { return array_interface_.num_cols; } + size_t NumRows() const { return array_interface_.Shape(0); } + size_t NumColumns() const { return array_interface_.Shape(1); } private: ArrayAdapterBatch batch_; - ArrayInterface array_interface_; + ArrayInterface<2> array_interface_; }; class CSRArrayAdapterBatch : public detail::NoMetaInfo { - ArrayInterface indptr_; - ArrayInterface indices_; - ArrayInterface values_; + ArrayInterface<1> indptr_; + ArrayInterface<1> indices_; + ArrayInterface<1> values_; bst_feature_t n_features_; class Line { - ArrayInterface indices_; - ArrayInterface values_; + ArrayInterface<1> indices_; + ArrayInterface<1> values_; size_t ridx_; size_t offset_; public: - Line(ArrayInterface indices, ArrayInterface values, size_t ridx, + Line(ArrayInterface<1> indices, ArrayInterface<1> values, size_t ridx, size_t offset) : indices_{std::move(indices)}, values_{std::move(values)}, ridx_{ridx}, offset_{offset} {} COOTuple GetElement(size_t idx) const { - return {ridx_, indices_.GetElement(offset_ + idx, 0), - values_.GetElement(offset_ + idx, 0)}; + return {ridx_, indices_.operator()(offset_ + idx, 0), + values_(offset_ + idx, 0)}; } size_t Size() const { - return values_.num_rows * values_.num_cols; + return values_.Shape(0) * values_.Shape(1); } }; @@ -339,17 +339,20 @@ class CSRArrayAdapterBatch : public detail::NoMetaInfo { public: CSRArrayAdapterBatch() = default; - CSRArrayAdapterBatch(ArrayInterface indptr, ArrayInterface indices, - ArrayInterface values, bst_feature_t n_features) - : indptr_{std::move(indptr)}, indices_{std::move(indices)}, - values_{std::move(values)}, n_features_{n_features} { - indptr_.AsColumnVector(); - values_.AsColumnVector(); - indices_.AsColumnVector(); + CSRArrayAdapterBatch(ArrayInterface<1> indptr, ArrayInterface<1> indices, + ArrayInterface<1> values, bst_feature_t n_features) + : indptr_{std::move(indptr)}, + indices_{std::move(indices)}, + values_{std::move(values)}, + n_features_{n_features} { + LOG(FATAL) << ""; + // indptr_.AsColumnVector(); + // values_.AsColumnVector(); + // indices_.AsColumnVector(); } size_t NumRows() const { - size_t size = indptr_.num_rows * indptr_.num_cols; + size_t size = indptr_.Shape(0) * indptr_.Shape(1); size = size == 0 ? 0 : size - 1; return size; } @@ -357,17 +360,18 @@ class CSRArrayAdapterBatch : public detail::NoMetaInfo { size_t Size() const { return this->NumRows(); } Line const GetLine(size_t idx) const { - auto begin_offset = indptr_.GetElement(idx, 0); - auto end_offset = indptr_.GetElement(idx + 1, 0); + auto begin_offset = indptr_.operator()(idx, 0); + auto end_offset = indptr_.operator()(idx + 1, 0); auto indices = indices_; auto values = values_; + LOG(FATAL) << ""; - values.num_cols = end_offset - begin_offset; - values.num_rows = 1; + // values.shape[0] = end_offset - begin_offset; + // indices.shape[0] - indices.num_cols = values.num_cols; - indices.num_rows = values.num_rows; + // indices.num_cols = values.Shape(1); + // indices.num_rows = values.Shape(0); return Line{indices, values, idx, begin_offset}; } @@ -391,7 +395,7 @@ class CSRArrayAdapter : public detail::SingleBatchDataIter return batch_; } size_t NumRows() const { - size_t size = indptr_.num_cols * indptr_.num_rows; + size_t size = indptr_.Shape(1) * indptr_.Shape(0); size = size == 0 ? 0 : size - 1; return size; } @@ -399,9 +403,9 @@ class CSRArrayAdapter : public detail::SingleBatchDataIter private: CSRArrayAdapterBatch batch_; - ArrayInterface indptr_; - ArrayInterface indices_; - ArrayInterface values_; + ArrayInterface<1> indptr_; + ArrayInterface<1> indices_; + ArrayInterface<1> values_; size_t num_cols_; }; diff --git a/src/data/array_interface.h b/src/data/array_interface.h index 6524f4512407..5ffe1ba421c8 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -13,24 +13,23 @@ #include #include +#include "../common/bitfield.h" +#include "../common/common.h" #include "xgboost/base.h" #include "xgboost/data.h" #include "xgboost/json.h" +#include "xgboost/linalg.h" #include "xgboost/logging.h" #include "xgboost/span.h" -#include "../common/bitfield.h" -#include "../common/common.h" namespace xgboost { // Common errors in parsing columnar format. struct ArrayInterfaceErrors { - static char const* Contigious() { - return "Memory should be contigious."; - } - static char const* TypestrFormat() { + static char const *Contigious() { return "Memory should be contigious."; } + static char const *TypestrFormat() { return "`typestr' should be of format ."; } - static char const* Dimension(int32_t d) { + static char const *Dimension(int32_t d) { static std::string str; str.clear(); str += "Only "; @@ -38,11 +37,11 @@ struct ArrayInterfaceErrors { str += " dimensional array is valid."; return str.c_str(); } - static char const* Version() { + static char const *Version() { return "Only version <= 3 of " "`__cuda_array_interface__/__array_interface__' are supported."; } - static char const* OfType(std::string const& type) { + static char const *OfType(std::string const &type) { static std::string str; str.clear(); str += " should be of "; @@ -96,28 +95,26 @@ struct ArrayInterfaceErrors { // object and turn it into an array (for cupy and numba). class ArrayInterfaceHandler { public: + enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 }; + template static constexpr char TypeChar() { - return - (std::is_floating_point::value ? 'f' : - (std::is_integral::value ? - (std::is_signed::value ? 'i' : 'u') : '\0')); + return (std::is_floating_point::value + ? 'f' + : (std::is_integral::value ? (std::is_signed::value ? 'i' : 'u') : '\0')); } template - static PtrType GetPtrFromArrayData(std::map const& obj) { + static PtrType GetPtrFromArrayData(std::map const &obj) { if (obj.find("data") == obj.cend()) { LOG(FATAL) << "Empty data passed in."; } - auto p_data = reinterpret_cast(static_cast( - get( - get( - obj.at("data")) - .at(0)))); + auto p_data = reinterpret_cast( + static_cast(get(get(obj.at("data")).at(0)))); return p_data; } - static void Validate(std::map const& array) { + static void Validate(std::map const &array) { auto version_it = array.find("version"); if (version_it == array.cend()) { LOG(FATAL) << "Missing `version' field for array interface"; @@ -149,12 +146,12 @@ class ArrayInterfaceHandler { // Mask object is also an array interface, but with different requirements. static size_t ExtractMask(std::map const &column, common::Span *p_out) { - auto& s_mask = *p_out; + auto &s_mask = *p_out; if (column.find("mask") != column.cend()) { - auto const& j_mask = get(column.at("mask")); + auto const &j_mask = get(column.at("mask")); Validate(j_mask); - auto p_mask = GetPtrFromArrayData(j_mask); + auto p_mask = GetPtrFromArrayData(j_mask); auto j_shape = get(j_mask.at("shape")); CHECK_EQ(j_shape.size(), 1) << ArrayInterfaceErrors::Dimension(1); @@ -186,7 +183,7 @@ class ArrayInterfaceHandler { if (j_mask.find("strides") != j_mask.cend()) { auto strides = get(column.at("strides")); - CHECK_EQ(strides.size(), 1) << ArrayInterfaceErrors::Dimension(1); + CHECK_EQ(strides.size(), 1) << ArrayInterfaceErrors::Dimension(1); CHECK_EQ(get(strides.at(0)), type_length) << ArrayInterfaceErrors::Contigious(); } @@ -196,51 +193,63 @@ class ArrayInterfaceHandler { return 0; } - static std::pair ExtractShape( - std::map const& column) { - auto j_shape = get(column.at("shape")); - auto typestr = get(column.at("typestr")); - if (j_shape.size() == 1) { - return {static_cast(get(j_shape.at(0))), 1}; - } else { - CHECK_EQ(j_shape.size(), 2) << "Only 1-D and 2-D arrays are supported."; - return {static_cast(get(j_shape.at(0))), - static_cast(get(j_shape.at(1)))}; + template + static void ExtractShape(std::map const &array, size_t (&out_shape)[D]) { + auto const &j_shape = get(array.at("shape")); + + size_t k = 0; + for (auto const& jval : j_shape) { + auto const &v = get(jval); + // Remove all empty dimension. + if (v == 0 || v == 1) { + continue; + } + CHECK_LT(k, D) << ArrayInterfaceErrors::Dimension(D); + out_shape[k] = v; + k++; + } + for (size_t i = j_shape.size(); i < D; ++i) { + out_shape[i] = 1; } } - static void ExtractStride(std::map const &column, - size_t *stride_r, size_t *stride_c, size_t rows, - size_t cols, size_t itemsize) { - auto strides_it = column.find("strides"); - if (strides_it == column.cend() || IsA(strides_it->second)) { - // default strides - *stride_r = cols; - *stride_c = 1; - } else { - // strides specified by the array interface - auto const &j_strides = get(strides_it->second); - CHECK_LE(j_strides.size(), 2) << ArrayInterfaceErrors::Dimension(2); - *stride_r = get(j_strides[0]) / itemsize; - size_t n = 1; - if (j_strides.size() == 2) { - n = get(j_strides[1]) / itemsize; - } - *stride_c = n; + template + static bool ExtractStride(std::map const &array, size_t itemsize, + size_t (&shape)[D], size_t (&stride)[D]) { + auto strides_it = array.find("strides"); + if (strides_it == array.cend() || IsA(strides_it->second)) { + // No stride is provided, we can calculate it from shape. + linalg::detail::CalcStride(shape, stride); + return true; } - auto valid = rows * (*stride_r) + cols * (*stride_c) >= (rows * cols); - CHECK(valid) << "Invalid strides in array." - << " strides: (" << (*stride_r) << "," << (*stride_c) - << "), shape: (" << rows << ", " << cols << ")"; + auto const &j_strides = get(strides_it->second); + auto const &j_shape = get(array.at("shape")); + CHECK_EQ(j_strides.size(), j_shape.size()) << "stride and shape don't match."; + + size_t k = 0; + for (size_t i = 0; i < j_strides.size(); ++i) { + auto const &stride_i = get(j_strides[i]); + auto const &shape_i = get(j_shape[i]); + // Remove empty dimension + if (shape_i == 0 || shape_i == 1) { + continue; + } + CHECK_LT(k, D) << ArrayInterfaceErrors::Dimension(D); + stride[k] = stride_i; + k++; + } + for (size_t i = j_strides.size(); i < D; ++i) { + stride[i] = 1; + } + return false; } - static void* ExtractData(std::map const &column, - std::pair shape) { - Validate(column); - void* p_data = ArrayInterfaceHandler::GetPtrFromArrayData(column); + static void *ExtractData(std::map const &array, size_t size) { + Validate(array); + void *p_data = ArrayInterfaceHandler::GetPtrFromArrayData(array); if (!p_data) { - CHECK_EQ(shape.first * shape.second, 0) << "Empty data with non-zero shape."; + CHECK_EQ(size, 0) << "Empty data with non-zero shape."; } return p_data; } @@ -249,42 +258,66 @@ class ArrayInterfaceHandler { }; #if !defined(XGBOOST_USE_CUDA) -inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { - common::AssertGPUSupport(); -} +inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { common::AssertGPUSupport(); } #endif // !defined(XGBOOST_USE_CUDA) // A view over __array_interface__ +/** + * \brief A type erased view over __array_interface__ protocol defined by numpy + * + * numpy. + * + * \tparam D The number of maximum dimension. + + * User input array must have dim <= D for all non-trivial dimensions. During + * construction, the ctor can automatically remove those trivial dimensions. + * + * \tparam allow_mask Whether masked array is accepted. + * + * Currently this only supported for 1-dim vector, which is used by cuDF column + * (apache arrow format). For general masked array, as the time of writting, only + * numpy has the proper support even though it's in the __cuda_array_interface__ + * protocol defined by numba. + */ +template class ArrayInterface { - void Initialize(std::map const &array, - bool allow_mask = true) { + static_assert(D > 0, "Invalid dimension for array interface."); + + /** + * \brief Initialize the object, by extracting shape, stride and type. + * + * The function also perform some basic validation for input array. Lastly it will + * also remove trivial dimensions like converting a matrix with shape (n_samples, 1) + * to a vector of size n_samples. For for inputs like weights, this should be a 1 + * dimension column vector even though user might provide a matrix. + */ + void Initialize(std::map const &array) { ArrayInterfaceHandler::Validate(array); + auto typestr = get(array.at("typestr")); this->AssignType(StringView{typestr}); + ArrayInterfaceHandler::ExtractShape(array, shape); + size_t itemsize = typestr[2] - '0'; + is_contigious = ArrayInterfaceHandler::ExtractStride(array, itemsize, shape, strides); + n = linalg::detail::CalcSize(shape); - std::tie(num_rows, num_cols) = ArrayInterfaceHandler::ExtractShape(array); - data = ArrayInterfaceHandler::ExtractData( - array, std::make_pair(num_rows, num_cols)); + data = ArrayInterfaceHandler::ExtractData(array, n); if (allow_mask) { + CHECK(D == 1) << "Masked array is not supported."; common::Span s_mask; size_t n_bits = ArrayInterfaceHandler::ExtractMask(array, &s_mask); valid = RBitField8(s_mask); if (s_mask.data()) { - CHECK_EQ(n_bits, num_rows) - << "Shape of bit mask doesn't match data shape. " - << "XGBoost doesn't support internal broadcasting."; + CHECK_EQ(n_bits, n) << "Shape of bit mask doesn't match data shape. " + << "XGBoost doesn't support internal broadcasting."; } } else { - CHECK(array.find("mask") == array.cend()) - << "Masked array is not yet supported."; + CHECK(array.find("mask") == array.cend()) << "Masked array is not yet supported."; } - ArrayInterfaceHandler::ExtractStride(array, &stride_row, &stride_col, - num_rows, num_cols, typestr[2] - '0'); - auto stream_it = array.find("stream"); if (stream_it != array.cend() && !IsA(stream_it->second)) { int64_t stream = get(stream_it->second); @@ -292,151 +325,123 @@ class ArrayInterface { } } - public: - enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 }; - public: ArrayInterface() = default; - explicit ArrayInterface(std::string const &str, bool allow_mask = true) - : ArrayInterface{StringView{str.c_str(), str.size()}, allow_mask} {} - - explicit ArrayInterface(std::map const &column, - bool allow_mask = true) { - this->Initialize(column, allow_mask); + explicit ArrayInterface(std::map const &array) { + this->Initialize(array); } - explicit ArrayInterface(StringView str, bool allow_mask = true) { + explicit ArrayInterface(std::string const &str) + : ArrayInterface{StringView{str}} {} + + explicit ArrayInterface(StringView str) { auto jinterface = Json::Load(str); if (IsA(jinterface)) { - this->Initialize(get(jinterface), allow_mask); + this->Initialize(get(jinterface)); return; } if (IsA(jinterface)) { CHECK_EQ(get(jinterface).size(), 1) << "Column: " << ArrayInterfaceErrors::Dimension(1); - this->Initialize(get(get(jinterface)[0]), allow_mask); + this->Initialize(get(get(jinterface)[0])); return; } } - void AsColumnVector() { - CHECK(num_rows == 1 || num_cols == 1) << "Array should be a vector instead of matrix."; - num_rows = std::max(num_rows, static_cast(num_cols)); - num_cols = 1; - - stride_row = std::max(stride_row, stride_col); - stride_col = 1; - } - void AssignType(StringView typestr) { - if (typestr.size() == 4 && typestr[1] == 'f' && typestr[2] == '1' && - typestr[3] == '6') { - type = kF16; + using T = ArrayInterfaceHandler::Type; + if (typestr.size() == 4 && typestr[1] == 'f' && typestr[2] == '1' && typestr[3] == '6') { + type = T::kF16; CHECK(sizeof(long double) == 16) << "128-bit floating point is not supported on current platform."; } else if (typestr[1] == 'f' && typestr[2] == '4') { - type = kF4; + type = T::kF4; } else if (typestr[1] == 'f' && typestr[2] == '8') { - type = kF8; + type = T::kF8; } else if (typestr[1] == 'i' && typestr[2] == '1') { - type = kI1; + type = T::kI1; } else if (typestr[1] == 'i' && typestr[2] == '2') { - type = kI2; + type = T::kI2; } else if (typestr[1] == 'i' && typestr[2] == '4') { - type = kI4; + type = T::kI4; } else if (typestr[1] == 'i' && typestr[2] == '8') { - type = kI8; + type = T::kI8; } else if (typestr[1] == 'u' && typestr[2] == '1') { - type = kU1; + type = T::kU1; } else if (typestr[1] == 'u' && typestr[2] == '2') { - type = kU2; + type = T::kU2; } else if (typestr[1] == 'u' && typestr[2] == '4') { - type = kU4; + type = T::kU4; } else if (typestr[1] == 'u' && typestr[2] == '8') { - type = kU8; + type = T::kU8; } else { LOG(FATAL) << ArrayInterfaceErrors::UnSupportedType(typestr); return; } } + XGBOOST_DEVICE size_t Shape(size_t i) const { return shape[i]; } + XGBOOST_DEVICE size_t Stride(size_t i) const { return strides[i]; } + template - XGBOOST_HOST_DEV_INLINE decltype(auto) DispatchCall(Fn func) const { + XGBOOST_HOST_DEV_INLINE constexpr decltype(auto) DispatchCall(Fn func) const { + using T = ArrayInterfaceHandler::Type; switch (type) { - case kF4: - return func(reinterpret_cast(data)); - case kF8: - return func(reinterpret_cast(data)); + case T::kF4: + return func(reinterpret_cast(data)); + case T::kF8: + return func(reinterpret_cast(data)); #ifdef __CUDA_ARCH__ - case kF16: { - // CUDA device code doesn't support long double. - SPAN_CHECK(false); - return func(reinterpret_cast(data)); - } + case T::kF16: { + // CUDA device code doesn't support long double. + SPAN_CHECK(false); + return func(reinterpret_cast(data)); + } #else - case kF16: - return func(reinterpret_cast(data)); + case T::kF16: + return func(reinterpret_cast(data)); #endif - case kI1: - return func(reinterpret_cast(data)); - case kI2: - return func(reinterpret_cast(data)); - case kI4: - return func(reinterpret_cast(data)); - case kI8: - return func(reinterpret_cast(data)); - case kU1: - return func(reinterpret_cast(data)); - case kU2: - return func(reinterpret_cast(data)); - case kU4: - return func(reinterpret_cast(data)); - case kU8: - return func(reinterpret_cast(data)); + case T::kI1: + return func(reinterpret_cast(data)); + case T::kI2: + return func(reinterpret_cast(data)); + case T::kI4: + return func(reinterpret_cast(data)); + case T::kI8: + return func(reinterpret_cast(data)); + case T::kU1: + return func(reinterpret_cast(data)); + case T::kU2: + return func(reinterpret_cast(data)); + case T::kU4: + return func(reinterpret_cast(data)); + case T::kU8: + return func(reinterpret_cast(data)); } SPAN_CHECK(false); return func(reinterpret_cast(data)); } - XGBOOST_DEVICE size_t ElementSize() { - return this->DispatchCall([](auto* p_values) { - return sizeof(std::remove_pointer_t); - }); + XGBOOST_DEVICE size_t constexpr ElementSize() { + return this->DispatchCall( + [](auto *p_values) { return sizeof(std::remove_pointer_t); }); } - template - XGBOOST_DEVICE T GetElement(size_t r, size_t c) const { - return this->DispatchCall( - [=](auto *p_values) -> T { return p_values[stride_row * r + stride_col * c]; }); + template + XGBOOST_DEVICE T operator()(Index &&...index) const { + return this->DispatchCall([=](auto const *p_values) -> T { + size_t offset = linalg::detail::Offset(strides, 0ul, 0ul, index...); + return static_cast(p_values[offset]); + }); } RBitField8 valid; - bst_row_t num_rows; - bst_feature_t num_cols; - size_t stride_row{0}; - size_t stride_col{0}; - void* data; - Type type; + size_t strides[D]{0}; + size_t shape[D]{0}; + void *data; + size_t n; + bool is_contigious {false}; + ArrayInterfaceHandler::Type type; }; - -template std::string MakeArrayInterface(T const *data, size_t n) { - Json arr{Object{}}; - arr["data"] = Array(std::vector{ - Json{Integer{reinterpret_cast(data)}}, Json{Boolean{false}}}); - arr["shape"] = Array{std::vector{Json{Integer{n}}, Json{Integer{1}}}}; - std::string typestr; - if (DMLC_LITTLE_ENDIAN) { - typestr.push_back('<'); - } else { - typestr.push_back('>'); - } - typestr.push_back(ArrayInterfaceHandler::TypeChar()); - typestr += std::to_string(sizeof(T)); - arr["typestr"] = typestr; - arr["version"] = 3; - std::string str; - Json::Dump(arr, &str); - return str; -} } // namespace xgboost #endif // XGBOOST_DATA_ARRAY_INTERFACE_H_ diff --git a/src/data/data.cc b/src/data/data.cc index 91526429e3b8..b7d72ddc2539 100644 --- a/src/data/data.cc +++ b/src/data/data.cc @@ -1,5 +1,5 @@ /*! - * Copyright 2015-2020 by Contributors + * Copyright 2015-2021 by Contributors * \file data.cc */ #include @@ -12,6 +12,8 @@ #include "xgboost/logging.h" #include "xgboost/version_config.h" #include "xgboost/learner.h" +#include "xgboost/string_view.h" + #include "sparse_page_writer.h" #include "simple_dmatrix.h" @@ -65,10 +67,22 @@ void SaveVectorField(dmlc::Stream* strm, const std::string& name, SaveVectorField(strm, name, type, shape, field.ConstHostVector()); } +template +void SaveTensorField(dmlc::Stream* strm, const std::string& name, xgboost::DataType type, + const xgboost::linalg::Tensor& field) { + strm->Write(name); + strm->Write(static_cast(type)); + strm->Write(false); // is_scalar=False + for (size_t i = 0; i < D; ++i) { + strm->Write(field.Shape(i)); + } + strm->Write(field.Data()->HostVector()); +} + template void LoadScalarField(dmlc::Stream* strm, const std::string& expected_name, xgboost::DataType expected_type, T* field) { - const std::string invalid {"MetaInfo: Invalid format. "}; + const std::string invalid{"MetaInfo: Invalid format for " + expected_name}; std::string name; xgboost::DataType type; bool is_scalar; @@ -90,7 +104,7 @@ void LoadScalarField(dmlc::Stream* strm, const std::string& expected_name, template void LoadVectorField(dmlc::Stream* strm, const std::string& expected_name, xgboost::DataType expected_type, std::vector* field) { - const std::string invalid {"MetaInfo: Invalid format. "}; + const std::string invalid{"MetaInfo: Invalid format for " + expected_name}; std::string name; xgboost::DataType type; bool is_scalar; @@ -123,6 +137,33 @@ void LoadVectorField(dmlc::Stream* strm, const std::string& expected_name, LoadVectorField(strm, expected_name, expected_type, &field->HostVector()); } +template +void LoadTensorField(dmlc::Stream* strm, std::string const& expected_name, + xgboost::DataType expected_type, xgboost::linalg::Tensor* p_out) { + const std::string invalid{"MetaInfo: Invalid format for " + expected_name}; + std::string name; + xgboost::DataType type; + bool is_scalar; + CHECK(strm->Read(&name)) << invalid; + CHECK_EQ(name, expected_name) << invalid << " Expected field: " << expected_name + << ", got: " << name; + uint8_t type_val; + CHECK(strm->Read(&type_val)) << invalid; + type = static_cast(type_val); + CHECK(type == expected_type) << invalid + << "Expected field of type: " << static_cast(expected_type) + << ", " + << "got field type: " << static_cast(type); + CHECK(strm->Read(&is_scalar)) << invalid; + CHECK(!is_scalar) << invalid << "Expected field " << expected_name + << " to be a vector; got a scalar"; + std::array shape; + for (size_t i = 0; i < D; ++i) { + CHECK(strm->Read(&(shape[i]))); + } + auto& field = p_out->Data()->HostVector(); + CHECK(strm->Read(&field)) << invalid; +} } // anonymous namespace namespace xgboost { @@ -135,7 +176,7 @@ void MetaInfo::Clear() { labels_.HostVector().clear(); group_ptr_.clear(); weights_.HostVector().clear(); - base_margin_.HostVector().clear(); + base_margin_ = decltype(base_margin_){}; } /* @@ -174,8 +215,7 @@ void MetaInfo::SaveBinary(dmlc::Stream *fo) const { {group_ptr_.size(), 1}, group_ptr_); ++field_cnt; SaveVectorField(fo, u8"weights", DataType::kFloat32, {weights_.Size(), 1}, weights_); ++field_cnt; - SaveVectorField(fo, u8"base_margin", DataType::kFloat32, - {base_margin_.Size(), 1}, base_margin_); ++field_cnt; + SaveTensorField(fo, u8"base_margin", DataType::kFloat32, base_margin_); ++field_cnt; SaveVectorField(fo, u8"labels_lower_bound", DataType::kFloat32, {labels_lower_bound_.Size(), 1}, labels_lower_bound_); ++field_cnt; SaveVectorField(fo, u8"labels_upper_bound", DataType::kFloat32, @@ -243,7 +283,7 @@ void MetaInfo::LoadBinary(dmlc::Stream *fi) { LoadVectorField(fi, u8"labels", DataType::kFloat32, &labels_); LoadVectorField(fi, u8"group_ptr", DataType::kUInt32, &group_ptr_); LoadVectorField(fi, u8"weights", DataType::kFloat32, &weights_); - LoadVectorField(fi, u8"base_margin", DataType::kFloat32, &base_margin_); + LoadTensorField(fi, u8"base_margin", DataType::kFloat32, &base_margin_); LoadVectorField(fi, u8"labels_lower_bound", DataType::kFloat32, &labels_lower_bound_); LoadVectorField(fi, u8"labels_upper_bound", DataType::kFloat32, &labels_upper_bound_); @@ -292,9 +332,10 @@ MetaInfo MetaInfo::Slice(common::Span ridxs) const { CHECK_EQ(this->base_margin_.Size() % this->num_row_, 0) << "Incorrect size of base margin vector."; size_t stride = this->base_margin_.Size() / this->num_row_; - out.base_margin_.HostVector() = Gather(this->base_margin_.HostVector(), ridxs, stride); + out.base_margin_.Data()->HostVector() = + Gather(this->base_margin_.Data()->HostVector(), ridxs, stride); } else { - out.base_margin_.HostVector() = Gather(this->base_margin_.HostVector(), ridxs); + out.base_margin_.Data()->HostVector() = Gather(this->base_margin_.Data()->HostVector(), ridxs); } out.feature_weigths.Resize(this->feature_weigths.Size()); @@ -364,7 +405,70 @@ void ValidateQueryGroup(std::vector const &group_ptr_) { auto cast_ptr = reinterpret_cast(old_ptr); proc; break; \ } \ default: LOG(FATAL) << "Unknown data type" << static_cast(dtype); \ - } \ + } + +namespace { +template +void CopyTensorInfoImpl(StringView interface_str, linalg::Tensor* p_out) { + ArrayInterface array{interface_str}; + p_out->ModifyInplace([&](HostDeviceVector* data, common::Span shape) { + data->Resize(array.n); + std::copy(array.shape, array.shape + D, shape.begin()); + auto t = p_out->View(GenericParameter::kCpuId); + common::ParallelFor(array.n, [&, span = data->HostSpan()](size_t i) { + span[i] = linalg::detail::Apply(array, linalg::detail::Unravel(i, shape)); + }); + }); +} +} // namespace + +void MetaInfo::SetInfo(StringView key, StringView interface_str) { + Json j_interface = Json::Load(interface_str); + bool is_cuda{false}; + if (IsA(j_interface)) { + auto const& array = get(j_interface); + CHECK_GE(array.size(), 0) << "Invalid " << key + << ", must have at least 1 column even if it's empty."; + auto const& first = get(array.front()); + is_cuda = first.find("stream") != first.cend(); + } else { + auto const& array = get(j_interface); + auto const& first = get(array.front()); + is_cuda = first.find("stream") != first.cend(); + } + + if (is_cuda) { + this->SetInfoFromCUDA(key.c_str(), std::string{interface_str.c_str()}); + } else { + this->SetInfoFromHost(key, interface_str); + } +} + +void MetaInfo::SetInfoFromHost(StringView key, StringView interface_str) { + if (key == "base_margin") { + CopyTensorInfoImpl<3>(interface_str, &this->base_margin_); + return; + } + + linalg::Tensor t; + CopyTensorInfoImpl<1>(interface_str, &t); + if (key == "label") { + this->labels_ = std::move(*t.Data()); + } else if (key == "weight") { + this->weights_ = std::move(*t.Data()); + } else if (key == "group") { + LOG(FATAL) << ""; + } else if (key == "qid") { + + } else if (key == "label_lower_bound") { + this->labels_lower_bound_ = std::move(*t.Data()); + } else if (key == "label_upper_bound") { + this->labels_upper_bound_ = std::move(*t.Data()); + } else if (key == "feature_weights") { + this->feature_weigths = std::move(*t.Data()); + } + LOG(FATAL) << "Unknown key for MetaInfo: " << key; +} void MetaInfo::SetInfo(const char* key, const void* dptr, DataType dtype, size_t num) { if (!std::strcmp(key, "label")) { @@ -386,7 +490,8 @@ void MetaInfo::SetInfo(const char* key, const void* dptr, DataType dtype, size_t }); CHECK(valid) << "Weights must be positive values."; } else if (!std::strcmp(key, "base_margin")) { - auto& base_margin = base_margin_.HostVector(); + // FIXME(jiamingy): Remove this as it doens't understand shape. + auto& base_margin = base_margin_.Data()->HostVector(); base_margin.resize(num); DISPATCH_CONST_PTR(dtype, dptr, cast_dptr, std::copy(cast_dptr, cast_dptr + num, base_margin.begin())); @@ -454,7 +559,7 @@ void MetaInfo::GetInfo(char const *key, bst_ulong *out_len, DataType dtype, } else if (!std::strcmp(key, "weight")) { vec = &this->weights_.HostVector(); } else if (!std::strcmp(key, "base_margin")) { - vec = &this->base_margin_.HostVector(); + vec = &this->base_margin_.Data()->HostVector(); } else if (!std::strcmp(key, "label_lower_bound")) { vec = &this->labels_lower_bound_.HostVector(); } else if (!std::strcmp(key, "label_upper_bound")) { @@ -543,8 +648,7 @@ void MetaInfo::Extend(MetaInfo const& that, bool accumulate_rows, bool check_col this->labels_upper_bound_.SetDevice(that.labels_upper_bound_.DeviceIdx()); this->labels_upper_bound_.Extend(that.labels_upper_bound_); - this->base_margin_.SetDevice(that.base_margin_.DeviceIdx()); - this->base_margin_.Extend(that.base_margin_); + linalg::Stack(&this->base_margin_, that.base_margin_); if (this->group_ptr_.size() == 0) { this->group_ptr_ = that.group_ptr_; @@ -627,7 +731,7 @@ void MetaInfo::Validate(int32_t device) const { if (base_margin_.Size() != 0) { CHECK_EQ(base_margin_.Size() % num_row_, 0) << "Size of base margin must be a multiple of number of rows."; - check_device(base_margin_); + check_device(*base_margin_.Data()); } } @@ -788,10 +892,10 @@ DMatrix* DMatrix::Load(const std::string& uri, LOG(CONSOLE) << info.group_ptr_.size() - 1 << " groups are loaded from " << fname << ".group"; } - if (MetaTryLoadFloatInfo - (fname + ".base_margin", &info.base_margin_.HostVector()) && !silent) { - LOG(CONSOLE) << info.base_margin_.Size() - << " base_margin are loaded from " << fname << ".base_margin"; + if (MetaTryLoadFloatInfo(fname + ".base_margin", &info.base_margin_.Data()->HostVector()) && + !silent) { + LOG(CONSOLE) << info.base_margin_.Size() << " base_margin are loaded from " << fname + << ".base_margin"; } if (MetaTryLoadFloatInfo (fname + ".weight", &info.weights_.HostVector()) && !silent) { diff --git a/src/data/data.cu b/src/data/data.cu index c4c7c503cb2e..25e66a523cc7 100644 --- a/src/data/data.cu +++ b/src/data/data.cu @@ -11,10 +11,72 @@ #include "../common/device_helpers.cuh" #include "device_adapter.cuh" #include "simple_dmatrix.h" +#include +#include namespace xgboost { +namespace { +// thrust::all_of tries to copy lambda function. +struct LabelsCheck { + __device__ bool operator()(float y) { return ::isnan(y) || ::isinf(y); } +}; +struct WeightsCheck { + __device__ bool operator()(float w) { return LabelsCheck{}(w) || w < 0; } // NOLINT +}; + +auto SetDeviceToPtr(void *ptr) { + cudaPointerAttributes attr; + dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); + int32_t ptr_device = attr.device; + dh::safe_cuda(cudaSetDevice(ptr_device)); + return ptr_device; +} + +// helper for type casting. +template +struct TypedGet { + ArrayInterface const& array; + + template + XGBOOST_DEV_INLINE auto operator()(Args&&... ind) { + return array.template operator()(std::forward(ind)...); + }; +}; + +template +void CopyTensorInfoImpl(ArrayInterface array, linalg::Tensor* p_out) { + if (array.n == 0) { + return; + } + auto ptr_device = SetDeviceToPtr(array.data); + + p_out->ModifyInplace([&](HostDeviceVector* data, common::Span shape) { + // set shape + std::copy(array.shape, array.shape + D, shape.begin()); + // set data + data->SetDevice(ptr_device); + data->Resize(array.n); + dh::LaunchN(array.n, [=, span = data->DeviceSpan()] __device__(size_t i) { + span[i] = linalg::detail::Apply( + [&array](auto&&... ind) { + // WAR: + // If the operator() is called directly, then there's this warning + // warning: parameter "ind" was declared but never referenced + return array.DispatchCall([&](auto const& p_values) -> T { + size_t offset = linalg::detail::Offset(array.strides, 0ul, 0ul, ind...); + return static_cast(p_values[offset]); + }); + }, + linalg::detail::Unravel(i, shape)); + }); + }); +} + +void CopyInfoImpl(ArrayInterface<1> column, HostDeviceVector* out) { + if (column.Shape(0) == 0) { + return; + } -void CopyInfoImpl(ArrayInterface column, HostDeviceVector* out) { auto SetDeviceToPtr = [](void* ptr) { cudaPointerAttributes attr; dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); @@ -26,67 +88,52 @@ void CopyInfoImpl(ArrayInterface column, HostDeviceVector* out) { }; auto ptr_device = SetDeviceToPtr(column.data); - if (column.num_rows == 0) { + if (column.Shape(0) == 0) { return; } out->SetDevice(ptr_device); + out->Resize(column.Shape(0)); - size_t size = column.num_rows * column.num_cols; + size_t size = column.n; CHECK_NE(size, 0); out->Resize(size); auto p_dst = thrust::device_pointer_cast(out->DevicePointer()); - dh::LaunchN(size, [=] __device__(size_t idx) { - size_t ridx = idx / column.num_cols; - size_t cidx = idx - (ridx * column.num_cols); - p_dst[idx] = column.GetElement(ridx, cidx); - }); -} - -namespace { -auto SetDeviceToPtr(void *ptr) { - cudaPointerAttributes attr; - dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); - int32_t ptr_device = attr.device; - dh::safe_cuda(cudaSetDevice(ptr_device)); - return ptr_device; + dh::LaunchN(column.Shape(0), + [=] __device__(size_t idx) { p_dst[idx] = column(idx, 0); }); } -} // anonymous namespace -void CopyGroupInfoImpl(ArrayInterface column, std::vector* out) { - CHECK(column.type != ArrayInterface::kF4 && column.type != ArrayInterface::kF8) +void CopyGroupInfoImpl(ArrayInterface<1> column, std::vector* out) { + CHECK(column.type != ArrayInterfaceHandler::kF4 && column.type != ArrayInterfaceHandler::kF8) << "Expected integer for group info."; auto ptr_device = SetDeviceToPtr(column.data); CHECK_EQ(ptr_device, dh::CurrentDevice()); - dh::TemporaryArray temp(column.num_rows); + dh::TemporaryArray temp(column.Shape(0)); auto d_tmp = temp.data(); - dh::LaunchN(column.num_rows, [=] __device__(size_t idx) { - d_tmp[idx] = column.GetElement(idx, 0); - }); - auto length = column.num_rows; + dh::LaunchN(column.Shape(0), + [=] __device__(size_t idx) { d_tmp[idx] = column.operator()(idx, 0); }); + auto length = column.Shape(0); out->resize(length + 1); out->at(0) = 0; thrust::copy(temp.data(), temp.data() + length, out->begin() + 1); std::partial_sum(out->begin(), out->end(), out->begin()); } -void CopyQidImpl(ArrayInterface array_interface, - std::vector *p_group_ptr) { +void CopyQidImpl(ArrayInterface<1> array_interface, std::vector* p_group_ptr) { auto &group_ptr_ = *p_group_ptr; - auto it = dh::MakeTransformIterator( - thrust::make_counting_iterator(0ul), - [array_interface] __device__(size_t i) { - return array_interface.GetElement(i, 0); - }); + auto it = dh::MakeTransformIterator(thrust::make_counting_iterator(0ul), + [array_interface] __device__(size_t i) { + return array_interface.operator()(i, 0); + }); dh::caching_device_vector flag(1); auto d_flag = dh::ToSpan(flag); auto d = SetDeviceToPtr(array_interface.data); dh::LaunchN(1, [=] __device__(size_t) { d_flag[0] = true; }); - dh::LaunchN(array_interface.num_rows - 1, [=] __device__(size_t i) { - if (array_interface.GetElement(i, 0) > - array_interface.GetElement(i + 1, 0)) { + dh::LaunchN(array_interface.Shape(0) - 1, [=] __device__(size_t i) { + if (array_interface.operator()(i, 0) > + array_interface.operator()(i + 1, 0)) { d_flag[0] = false; } }); @@ -95,16 +142,16 @@ void CopyQidImpl(ArrayInterface array_interface, cudaMemcpyDeviceToHost)); CHECK(non_dec) << "`qid` must be sorted in increasing order along with data."; size_t bytes = 0; - dh::caching_device_vector out(array_interface.num_rows); - dh::caching_device_vector cnt(array_interface.num_rows); + dh::caching_device_vector out(array_interface.Shape(0)); + dh::caching_device_vector cnt(array_interface.Shape(0)); HostDeviceVector d_num_runs_out(1, 0, d); cub::DeviceRunLengthEncode::Encode( nullptr, bytes, it, out.begin(), cnt.begin(), - d_num_runs_out.DevicePointer(), array_interface.num_rows); + d_num_runs_out.DevicePointer(), array_interface.Shape(0)); dh::caching_device_vector tmp(bytes); cub::DeviceRunLengthEncode::Encode( tmp.data().get(), bytes, it, out.begin(), cnt.begin(), - d_num_runs_out.DevicePointer(), array_interface.num_rows); + d_num_runs_out.DevicePointer(), array_interface.Shape(0)); auto h_num_runs_out = d_num_runs_out.HostSpan()[0]; group_ptr_.clear(); @@ -116,72 +163,60 @@ void CopyQidImpl(ArrayInterface array_interface, group_ptr_.begin() + 1); } -namespace { -// thrust::all_of tries to copy lambda function. -struct LabelsCheck { - __device__ bool operator()(float y) { return ::isnan(y) || ::isinf(y); } -}; -struct WeightsCheck { - __device__ bool operator()(float w) { return LabelsCheck{}(w) || w < 0; } // NOLINT -}; -} // anonymous namespace +template +void CheckArrayInterface(StringView key, ArrayInterface const& array) { + CHECK(!array.valid.Data()) << "Meta info " << key << " should be dense, found validity mask"; +} +} // namespace void ValidateQueryGroup(std::vector const &group_ptr_); -void MetaInfo::SetInfo(const char * c_key, std::string const& interface_str) { +void MetaInfo::SetInfoFromCUDA(const char* c_key, std::string const& interface_str) { Json j_interface = Json::Load({interface_str.c_str(), interface_str.size()}); - ArrayInterface array_interface(interface_str); - std::string key{c_key}; + auto const& j_arr = get(j_interface); + CHECK_EQ(j_arr.size(), 1) << "MetaInfo: " << c_key << ". " << ArrayInterfaceErrors::Dimension(1); - CHECK(!array_interface.valid.Data()) - << "Meta info " << key << " should be dense, found validity mask"; - if (array_interface.num_rows == 0) { - return; - } - - if (key == "base_margin") { - CopyInfoImpl(array_interface, &base_margin_); - return; - } + StringView key{c_key}; - CHECK(array_interface.num_cols == 1 || array_interface.num_rows == 1) - << "MetaInfo: " << c_key << " has invalid shape"; - if (!((array_interface.num_cols == 1 && array_interface.num_rows == 0) || - (array_interface.num_cols == 0 && array_interface.num_rows == 1))) { - // Not an empty column, transform it. - array_interface.AsColumnVector(); - } if (key == "label") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyInfoImpl(array_interface, &labels_); auto ptr = labels_.ConstDevicePointer(); - auto valid = thrust::none_of(thrust::device, ptr, ptr + labels_.Size(), - LabelsCheck{}); + auto valid = thrust::none_of(thrust::device, ptr, ptr + labels_.Size(), LabelsCheck{}); CHECK(valid) << "Label contains NaN, infinity or a value too large."; } else if (key == "weight") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyInfoImpl(array_interface, &weights_); auto ptr = weights_.ConstDevicePointer(); - auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(), - WeightsCheck{}); + auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(), WeightsCheck{}); CHECK(valid) << "Weights must be positive values."; + } else if (key == "base_margin") { + ArrayInterface<3> array_interface(StringView{interface_str}); + CopyTensorInfoImpl(array_interface, &base_margin_); } else if (key == "group") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyGroupInfoImpl(array_interface, &group_ptr_); ValidateQueryGroup(group_ptr_); return; } else if (key == "qid") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyQidImpl(array_interface, &group_ptr_); return; } else if (key == "label_lower_bound") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyInfoImpl(array_interface, &labels_lower_bound_); return; } else if (key == "label_upper_bound") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyInfoImpl(array_interface, &labels_upper_bound_); return; } else if (key == "feature_weights") { + auto array_interface{ArrayInterface<1>(StringView{interface_str})}; CopyInfoImpl(array_interface, &feature_weigths); auto d_feature_weights = feature_weigths.ConstDeviceSpan(); - auto valid = thrust::none_of( - thrust::device, d_feature_weights.data(), - d_feature_weights.data() + d_feature_weights.size(), WeightsCheck{}); + auto valid = + thrust::none_of(thrust::device, d_feature_weights.data(), + d_feature_weights.data() + d_feature_weights.size(), WeightsCheck{}); CHECK(valid) << "Feature weight must be greater than 0."; return; } else { diff --git a/src/data/device_adapter.cuh b/src/data/device_adapter.cuh index 628878f319f1..15474d904410 100644 --- a/src/data/device_adapter.cuh +++ b/src/data/device_adapter.cuh @@ -20,7 +20,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo { public: CudfAdapterBatch() = default; - CudfAdapterBatch(common::Span columns, size_t num_rows) + CudfAdapterBatch(common::Span> columns, size_t num_rows) : columns_(columns), num_rows_(num_rows) {} size_t Size() const { return num_rows_ * columns_.size(); } @@ -29,7 +29,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo { size_t row_idx = idx / columns_.size(); auto const& column = columns_[column_idx]; float value = column.valid.Data() == nullptr || column.valid.Check(row_idx) - ? column.GetElement(row_idx, 0) + ? column(row_idx, 0) : std::numeric_limits::quiet_NaN(); return {row_idx, column_idx, value}; } @@ -38,7 +38,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo { XGBOOST_DEVICE bst_row_t NumCols() const { return columns_.size(); } private: - common::Span columns_; + common::Span> columns_; size_t num_rows_; }; @@ -101,9 +101,9 @@ class CudfAdapter : public detail::SingleBatchDataIter { auto const& typestr = get(json_columns[0]["typestr"]); CHECK_EQ(typestr.size(), 3) << ArrayInterfaceErrors::TypestrFormat(); - std::vector columns; - auto first_column = ArrayInterface(get(json_columns[0])); - num_rows_ = first_column.num_rows; + std::vector> columns; + auto first_column = ArrayInterface<1>(get(json_columns[0])); + num_rows_ = first_column.Shape(0); if (num_rows_ == 0) { return; } @@ -112,13 +112,13 @@ class CudfAdapter : public detail::SingleBatchDataIter { CHECK_NE(device_idx_, -1); dh::safe_cuda(cudaSetDevice(device_idx_)); for (auto& json_col : json_columns) { - auto column = ArrayInterface(get(json_col)); + auto column = ArrayInterface<1>(get(json_col)); columns.push_back(column); - CHECK_EQ(column.num_cols, 1); - num_rows_ = std::max(num_rows_, size_t(column.num_rows)); + CHECK_EQ(column.Shape(1), 1); + num_rows_ = std::max(num_rows_, size_t(column.Shape(0))); CHECK_EQ(device_idx_, dh::CudaGetPointerDevice(column.data)) << "All columns should use the same device."; - CHECK_EQ(num_rows_, column.num_rows) + CHECK_EQ(num_rows_, column.Shape(0)) << "All columns should have same number of rows."; } columns_ = columns; @@ -135,7 +135,7 @@ class CudfAdapter : public detail::SingleBatchDataIter { private: CudfAdapterBatch batch_; - dh::device_vector columns_; + dh::device_vector> columns_; size_t num_rows_{0}; int device_idx_; }; @@ -143,23 +143,23 @@ class CudfAdapter : public detail::SingleBatchDataIter { class CupyAdapterBatch : public detail::NoMetaInfo { public: CupyAdapterBatch() = default; - explicit CupyAdapterBatch(ArrayInterface array_interface) + explicit CupyAdapterBatch(ArrayInterface<2> array_interface) : array_interface_(std::move(array_interface)) {} size_t Size() const { - return array_interface_.num_rows * array_interface_.num_cols; + return array_interface_.Shape(0) * array_interface_.Shape(1); } __device__ COOTuple GetElement(size_t idx) const { - size_t column_idx = idx % array_interface_.num_cols; - size_t row_idx = idx / array_interface_.num_cols; - float value = array_interface_.GetElement(row_idx, column_idx); + size_t column_idx = idx % array_interface_.Shape(1); + size_t row_idx = idx / array_interface_.Shape(1); + float value = array_interface_(row_idx, column_idx); return {row_idx, column_idx, value}; } - XGBOOST_DEVICE bst_row_t NumRows() const { return array_interface_.num_rows; } - XGBOOST_DEVICE bst_row_t NumCols() const { return array_interface_.num_cols; } + XGBOOST_DEVICE bst_row_t NumRows() const { return array_interface_.Shape(0); } + XGBOOST_DEVICE bst_row_t NumCols() const { return array_interface_.Shape(1); } private: - ArrayInterface array_interface_; + ArrayInterface<2> array_interface_; }; class CupyAdapter : public detail::SingleBatchDataIter { @@ -167,9 +167,9 @@ class CupyAdapter : public detail::SingleBatchDataIter { explicit CupyAdapter(std::string cuda_interface_str) { Json json_array_interface = Json::Load({cuda_interface_str.c_str(), cuda_interface_str.size()}); - array_interface_ = ArrayInterface(get(json_array_interface), false); + array_interface_ = ArrayInterface<2>(get(json_array_interface)); batch_ = CupyAdapterBatch(array_interface_); - if (array_interface_.num_rows == 0) { + if (array_interface_.Shape(0) == 0) { return; } device_idx_ = dh::CudaGetPointerDevice(array_interface_.data); @@ -177,12 +177,12 @@ class CupyAdapter : public detail::SingleBatchDataIter { } const CupyAdapterBatch& Value() const override { return batch_; } - size_t NumRows() const { return array_interface_.num_rows; } - size_t NumColumns() const { return array_interface_.num_cols; } + size_t NumRows() const { return array_interface_.Shape(0); } + size_t NumColumns() const { return array_interface_.Shape(1); } int32_t DeviceIdx() const { return device_idx_; } private: - ArrayInterface array_interface_; + ArrayInterface<2> array_interface_; CupyAdapterBatch batch_; int32_t device_idx_ {-1}; }; diff --git a/src/data/file_iterator.h b/src/data/file_iterator.h index 6d6adb62b008..63060f4f4f0b 100644 --- a/src/data/file_iterator.h +++ b/src/data/file_iterator.h @@ -12,6 +12,7 @@ #include "dmlc/data.h" #include "xgboost/c_api.h" #include "xgboost/json.h" +#include "xgboost/linalg.h" #include "array_interface.h" namespace xgboost { @@ -58,16 +59,14 @@ class FileIterator { CHECK(parser_); if (parser_->Next()) { row_block_ = parser_->Value(); + using linalg::MakeVec; - indptr_ = MakeArrayInterface(row_block_.offset, row_block_.size + 1); - values_ = MakeArrayInterface(row_block_.value, - row_block_.offset[row_block_.size]); - indices_ = MakeArrayInterface(row_block_.index, - row_block_.offset[row_block_.size]); + indptr_ = MakeVec(row_block_.offset, row_block_.size + 1).ArrayInterface(); + values_ = MakeVec(row_block_.value, row_block_.offset[row_block_.size]).ArrayInterface(); + indices_ = MakeVec(row_block_.index, row_block_.offset[row_block_.size]).ArrayInterface(); - size_t n_columns = *std::max_element( - row_block_.index, - row_block_.index + row_block_.offset[row_block_.size]); + size_t n_columns = *std::max_element(row_block_.index, + row_block_.index + row_block_.offset[row_block_.size]); // dmlc parser converts 1-based indexing back to 0-based indexing so we can ignore // this condition and just add 1 to n_columns n_columns += 1; diff --git a/src/data/simple_dmatrix.cc b/src/data/simple_dmatrix.cc index 44a8a3f8fe7c..e83559d3958a 100644 --- a/src/data/simple_dmatrix.cc +++ b/src/data/simple_dmatrix.cc @@ -137,9 +137,10 @@ SimpleDMatrix::SimpleDMatrix(AdapterT* adapter, float missing, int nthread) { batch.Weights() + batch.Size()); } if (batch.BaseMargin() != nullptr) { - auto& base_margin = info_.base_margin_.HostVector(); - base_margin.insert(base_margin.end(), batch.BaseMargin(), - batch.BaseMargin() + batch.Size()); + info_.base_margin_ = linalg::Tensor{batch.BaseMargin(), + batch.BaseMargin() + batch.Size(), + {batch.Size()}, + GenericParameter::kCpuId}; } if (batch.Qid() != nullptr) { qids.insert(qids.end(), batch.Qid(), batch.Qid() + batch.Size()); diff --git a/src/gbm/gblinear.cc b/src/gbm/gblinear.cc index e5f2a457214c..e142ef5758ba 100644 --- a/src/gbm/gblinear.cc +++ b/src/gbm/gblinear.cc @@ -178,7 +178,7 @@ class GBLinear : public GradientBooster { unsigned layer_begin, unsigned layer_end, bool, int, unsigned) override { model_.LazyInitModel(); LinearCheckLayer(layer_begin, layer_end); - const auto& base_margin = p_fmat->Info().base_margin_.ConstHostVector(); + auto base_margin = p_fmat->Info().base_margin_.View(GenericParameter::kCpuId); const int ngroup = model_.learner_model_param->num_output_group; const size_t ncolumns = model_.learner_model_param->num_feature + 1; // allocate space for (#features + bias) times #groups times #rows @@ -203,9 +203,9 @@ class GBLinear : public GradientBooster { p_contribs[ins.index] = ins.fvalue * model_[ins.index][gid]; } // add base margin to BIAS - p_contribs[ncolumns - 1] = model_.Bias()[gid] + - ((base_margin.size() != 0) ? base_margin[row_idx * ngroup + gid] : - learner_model_param_->base_score); + p_contribs[ncolumns - 1] = + model_.Bias()[gid] + ((base_margin.Size() != 0) ? base_margin(row_idx, gid) + : learner_model_param_->base_score); } }); } @@ -250,8 +250,10 @@ class GBLinear : public GradientBooster { {learner_model_param_->num_feature, n_groups}, GenericParameter::kCpuId}; for (size_t i = 0; i < learner_model_param_->num_feature; ++i) { + auto feature_score = scores.Slice(i, linalg::All()); + CHECK_EQ(feature_score.Shape().size(), 1); for (bst_group_t g = 0; g < n_groups; ++g) { - scores(i, g) = model_[i][g]; + feature_score(g) = model_[i][g]; } } } @@ -270,7 +272,7 @@ class GBLinear : public GradientBooster { monitor_.Start("PredictBatchInternal"); model_.LazyInitModel(); std::vector &preds = *out_preds; - const auto& base_margin = p_fmat->Info().base_margin_.ConstHostVector(); + auto base_margin = p_fmat->Info().base_margin_.View(GenericParameter::kCpuId); // start collecting the prediction const int ngroup = model_.learner_model_param->num_output_group; preds.resize(p_fmat->Info().num_row_ * ngroup); @@ -280,16 +282,15 @@ class GBLinear : public GradientBooster { // k is number of group // parallel over local batch const auto nsize = static_cast(batch.Size()); - if (base_margin.size() != 0) { - CHECK_EQ(base_margin.size(), nsize * ngroup); + if (base_margin.Size() != 0) { + CHECK_EQ(base_margin.Size(), nsize * ngroup); } common::ParallelFor(nsize, [&](omp_ulong i) { const size_t ridx = page.base_rowid + i; // loop over output groups for (int gid = 0; gid < ngroup; ++gid) { - bst_float margin = - (base_margin.size() != 0) ? - base_margin[ridx * ngroup + gid] : learner_model_param_->base_score; + float margin = (base_margin.Size() != 0) ? base_margin(ridx * ngroup + gid) + : learner_model_param_->base_score; this->Pred(batch[i], &preds[ridx * ngroup], gid, margin); } }); diff --git a/src/predictor/cpu_predictor.cc b/src/predictor/cpu_predictor.cc index d581f64a1d56..37f7d85e242c 100644 --- a/src/predictor/cpu_predictor.cc +++ b/src/predictor/cpu_predictor.cc @@ -287,18 +287,17 @@ class CPUPredictor : public Predictor { const gbm::GBTreeModel& model) const override { CHECK_NE(model.learner_model_param->num_output_group, 0); size_t n = model.learner_model_param->num_output_group * info.num_row_; - const auto& base_margin = info.base_margin_.HostVector(); + // const auto& base_margin = info.base_margin_.HostVector(); + auto base_margin = info.base_margin_.View(GenericParameter::kCpuId); out_preds->Resize(n); std::vector& out_preds_h = out_preds->HostVector(); - if (base_margin.empty()) { - std::fill(out_preds_h.begin(), out_preds_h.end(), - model.learner_model_param->base_score); + if (base_margin.Size() == 0) { + std::fill(out_preds_h.begin(), out_preds_h.end(), model.learner_model_param->base_score); } else { std::string expected{ "(" + std::to_string(info.num_row_) + ", " + std::to_string(model.learner_model_param->num_output_group) + ")"}; - CHECK_EQ(base_margin.size(), n) - << "Invalid shape of base_margin. Expected:" << expected; + CHECK_EQ(base_margin.Size(), n) << "Invalid shape of base_margin. Expected:" << expected; std::copy(base_margin.begin(), base_margin.end(), out_preds_h.begin()); } } @@ -456,7 +455,7 @@ class CPUPredictor : public Predictor { common::ParallelFor(bst_omp_uint(ntree_limit), [&](bst_omp_uint i) { FillNodeMeanValues(model.trees[i].get(), &(mean_values[i])); }); - const std::vector& base_margin = info.base_margin_.HostVector(); + auto base_margin = info.base_margin_.View(GenericParameter::kCpuId); // start collecting the contributions for (const auto &batch : p_fmat->GetBatches()) { auto page = batch.GetView(); @@ -496,8 +495,9 @@ class CPUPredictor : public Predictor { } feats.Drop(page[i]); // add base margin to BIAS - if (base_margin.size() != 0) { - p_contribs[ncolumns - 1] += base_margin[row_idx * ngroup + gid]; + if (base_margin.Size() != 0) { + CHECK_EQ(base_margin.Shape(1), ngroup); + p_contribs[ncolumns - 1] += base_margin(row_idx, gid); } else { p_contribs[ncolumns - 1] += model.learner_model_param->base_score; } diff --git a/src/predictor/gpu_predictor.cu b/src/predictor/gpu_predictor.cu index 51674237e973..38676e1db96c 100644 --- a/src/predictor/gpu_predictor.cu +++ b/src/predictor/gpu_predictor.cu @@ -854,8 +854,8 @@ class GPUPredictor : public xgboost::Predictor { dh::tend(phis)); } // Add the base margin term to last column - p_fmat->Info().base_margin_.SetDevice(generic_param_->gpu_id); - const auto margin = p_fmat->Info().base_margin_.ConstDeviceSpan(); + p_fmat->Info().base_margin_.Data()->SetDevice(generic_param_->gpu_id); + const auto margin = p_fmat->Info().base_margin_.Data()->ConstDeviceSpan(); float base_score = model.learner_model_param->base_score; dh::LaunchN( p_fmat->Info().num_row_ * model.learner_model_param->num_output_group, @@ -913,8 +913,8 @@ class GPUPredictor : public xgboost::Predictor { dh::tend(phis)); } // Add the base margin term to last column - p_fmat->Info().base_margin_.SetDevice(generic_param_->gpu_id); - const auto margin = p_fmat->Info().base_margin_.ConstDeviceSpan(); + p_fmat->Info().base_margin_.Data()->SetDevice(generic_param_->gpu_id); + const auto margin = p_fmat->Info().base_margin_.Data()->ConstDeviceSpan(); float base_score = model.learner_model_param->base_score; size_t n_features = model.learner_model_param->num_feature; dh::LaunchN( @@ -934,16 +934,12 @@ class GPUPredictor : public xgboost::Predictor { const gbm::GBTreeModel& model) const override { size_t n_classes = model.learner_model_param->num_output_group; size_t n = n_classes * info.num_row_; - const HostDeviceVector& base_margin = info.base_margin_; + const HostDeviceVector* base_margin = info.base_margin_.Data(); out_preds->SetDevice(generic_param_->gpu_id); out_preds->Resize(n); - if (base_margin.Size() != 0) { - std::string expected{ - "(" + std::to_string(info.num_row_) + ", " + - std::to_string(model.learner_model_param->num_output_group) + ")"}; - CHECK_EQ(base_margin.Size(), n) - << "Invalid shape of base_margin. Expected:" << expected; - out_preds->Copy(base_margin); + if (base_margin->Size() != 0) { + CHECK_EQ(base_margin->Size(), n); + out_preds->Copy(*base_margin); } else { out_preds->Fill(model.learner_model_param->base_score); } diff --git a/tests/cpp/common/test_linalg.cc b/tests/cpp/common/test_linalg.cc index e4ca2d86594b..a91bdb0b5e3e 100644 --- a/tests/cpp/common/test_linalg.cc +++ b/tests/cpp/common/test_linalg.cc @@ -1,4 +1,8 @@ +/*! + * Copyright 2021 by XGBoost Contributors + */ #include +#include #include #include @@ -6,6 +10,10 @@ namespace xgboost { namespace linalg { +namespace { +auto kCpuId = GenericParameter::kCpuId; +} + auto MakeMatrixFromTest(HostDeviceVector *storage, size_t n_rows, size_t n_cols) { storage->Resize(n_rows * n_cols); auto &h_storage = storage->HostVector(); @@ -16,16 +24,16 @@ auto MakeMatrixFromTest(HostDeviceVector *storage, size_t n_rows, size_t return m; } -TEST(Linalg, Matrix) { +TEST(Linalg, MatrixView) { size_t kRows = 31, kCols = 77; HostDeviceVector storage; auto m = MakeMatrixFromTest(&storage, kRows, kCols); - ASSERT_EQ(m.DeviceIdx(), GenericParameter::kCpuId); + ASSERT_EQ(m.DeviceIdx(), kCpuId); ASSERT_EQ(m(0, 0), 0); ASSERT_EQ(m(kRows - 1, kCols - 1), storage.Size() - 1); } -TEST(Linalg, Vector) { +TEST(Linalg, VectorView) { size_t kRows = 31, kCols = 77; HostDeviceVector storage; auto m = MakeMatrixFromTest(&storage, kRows, kCols); @@ -37,7 +45,7 @@ TEST(Linalg, Vector) { ASSERT_EQ(v(0), 3); } -TEST(Linalg, Tensor) { +TEST(Linalg, TensorView) { std::vector data(2 * 3 * 4, 0); std::iota(data.begin(), data.end(), 0); @@ -99,14 +107,50 @@ TEST(Linalg, Tensor) { } } +TEST(Linalg, Tensor) { + Tensor t{{2, 3, 4}, kCpuId}; + auto view = t.View(kCpuId); + + auto const &as_const = t; + auto k_view = as_const.View(kCpuId); + + size_t n = 2 * 3 * 4; + ASSERT_EQ(t.Size(), n); + ASSERT_TRUE(std::equal(k_view.cbegin(), k_view.cbegin(), view.begin())); +} + TEST(Linalg, Empty) { - auto t = TensorView{{}, {0, 3}, GenericParameter::kCpuId}; - for (int32_t i : {0, 1, 2}) { - auto s = t.Slice(All(), i); - ASSERT_EQ(s.Size(), 0); - ASSERT_EQ(s.Shape().size(), 1); - ASSERT_EQ(s.Shape(0), 0); + { + auto t = TensorView{{}, {0, 3}, kCpuId}; + for (int32_t i : {0, 1, 2}) { + auto s = t.Slice(All(), i); + ASSERT_EQ(s.Size(), 0); + ASSERT_EQ(s.Shape().size(), 1); + ASSERT_EQ(s.Shape(0), 0); + } + } + { + auto t = Tensor{{0, 3}, kCpuId}; + ASSERT_EQ(t.Size(), 0); + auto view = t.View(kCpuId); + + for (int32_t i : {0, 1, 2}) { + auto s = view.Slice(All(), i); + ASSERT_EQ(s.Size(), 0); + ASSERT_EQ(s.Shape().size(), 1); + ASSERT_EQ(s.Shape(0), 0); + } } } + +TEST(Linalg, ArrayInterface) { + auto cpu = kCpuId; + auto t = Tensor{{3, 3}, cpu}; + auto v = t.View(cpu); + std::iota(v.begin(), v.end(), 0); + auto arr = Json::Load(StringView{v.ArrayInterface()}); + ASSERT_EQ(get(arr["shape"][0]), 3); + ASSERT_EQ(get(arr["stride"][0]), 3 * sizeof(double)); +} } // namespace linalg } // namespace xgboost diff --git a/tests/cpp/data/test_adapter.cc b/tests/cpp/data/test_adapter.cc index ccb19de71a74..e2d0bd3769d9 100644 --- a/tests/cpp/data/test_adapter.cc +++ b/tests/cpp/data/test_adapter.cc @@ -41,9 +41,10 @@ TEST(Adapter, CSRArrayAdapter) { HostDeviceVector indices; size_t n_features = 100, n_samples = 10; RandomDataGenerator{n_samples, n_features, 0.5}.GenerateCSR(&values, &indptr, &indices); - auto indptr_arr = MakeArrayInterface(indptr.HostPointer(), indptr.Size()); - auto values_arr = MakeArrayInterface(values.HostPointer(), values.Size()); - auto indices_arr = MakeArrayInterface(indices.HostPointer(), indices.Size()); + using linalg::MakeVec; + auto indptr_arr = MakeVec(indptr.HostPointer(), indptr.Size()).ArrayInterface(); + auto values_arr = MakeVec(values.HostPointer(), values.Size()).ArrayInterface(); + auto indices_arr = MakeVec(indices.HostPointer(), indices.Size()).ArrayInterface(); auto adapter = data::CSRArrayAdapter( StringView{indptr_arr.c_str(), indptr_arr.size()}, StringView{values_arr.c_str(), values_arr.size()}, diff --git a/tests/cpp/data/test_array_interface.cc b/tests/cpp/data/test_array_interface.cc index 875858855ed5..d8adc4223753 100644 --- a/tests/cpp/data/test_array_interface.cc +++ b/tests/cpp/data/test_array_interface.cc @@ -11,21 +11,21 @@ TEST(ArrayInterface, Initialize) { size_t constexpr kRows = 10, kCols = 10; HostDeviceVector storage; auto array = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage); - auto arr_interface = ArrayInterface(array); - ASSERT_EQ(arr_interface.num_rows, kRows); - ASSERT_EQ(arr_interface.num_cols, kCols); + auto arr_interface = ArrayInterface<2>(StringView{array}); + ASSERT_EQ(arr_interface.Shape(0), kRows); + ASSERT_EQ(arr_interface.Shape(1), kCols); ASSERT_EQ(arr_interface.data, storage.ConstHostPointer()); ASSERT_EQ(arr_interface.ElementSize(), 4); - ASSERT_EQ(arr_interface.type, ArrayInterface::kF4); + ASSERT_EQ(arr_interface.type, ArrayInterfaceHandler::kF4); HostDeviceVector u64_storage(storage.Size()); std::string u64_arr_str; Json::Dump(GetArrayInterface(&u64_storage, kRows, kCols), &u64_arr_str); std::copy(storage.ConstHostVector().cbegin(), storage.ConstHostVector().cend(), u64_storage.HostSpan().begin()); - auto u64_arr = ArrayInterface{u64_arr_str}; + auto u64_arr = ArrayInterface<2>{u64_arr_str}; ASSERT_EQ(u64_arr.ElementSize(), 8); - ASSERT_EQ(u64_arr.type, ArrayInterface::kU8); + ASSERT_EQ(u64_arr.type, ArrayInterfaceHandler::kU8); } TEST(ArrayInterface, Error) { @@ -38,23 +38,22 @@ TEST(ArrayInterface, Error) { Json(Boolean(false))}; auto const& column_obj = get(column); - std::pair shape{kRows, kCols}; std::string typestr{"(1)); // missing data - EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape), + EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n), dmlc::Error); column["data"] = j_data; // missing typestr - EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape), + EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n), dmlc::Error); column["typestr"] = String(" storage; @@ -63,22 +62,41 @@ TEST(ArrayInterface, Error) { Json(Integer(reinterpret_cast(storage.ConstHostPointer()))), Json(Boolean(false))}; column["data"] = j_data; - EXPECT_NO_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape)); + EXPECT_NO_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n)); } TEST(ArrayInterface, GetElement) { size_t kRows = 4, kCols = 2; HostDeviceVector storage; auto intefrace_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage); - ArrayInterface array_interface{intefrace_str}; + ArrayInterface<2> array_interface{intefrace_str}; auto const& h_storage = storage.ConstHostVector(); for (size_t i = 0; i < kRows; ++i) { for (size_t j = 0; j < kCols; ++j) { - float v0 = array_interface.GetElement(i, j); + float v0 = array_interface(i, j); float v1 = h_storage.at(i * kCols + j); ASSERT_EQ(v0, v1); } } } + +TEST(ArrayInterface, TrivialDim) { + size_t kRows{1000}, kCols = 1; + HostDeviceVector storage; + auto interface_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage); + { + ArrayInterface<1> arr_i{interface_str}; + ASSERT_EQ(arr_i.n, kRows); + ASSERT_EQ(arr_i.Shape(0), kRows); + } + + std::swap(kRows, kCols); + interface_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage); + { + ArrayInterface<1> arr_i{interface_str}; + ASSERT_EQ(arr_i.n, kCols); + ASSERT_EQ(arr_i.Shape(0), kCols); + } +} } // namespace xgboost diff --git a/tests/cpp/data/test_array_interface.cu b/tests/cpp/data/test_array_interface.cu index 75923e74ba1a..db4f9eb1d701 100644 --- a/tests/cpp/data/test_array_interface.cu +++ b/tests/cpp/data/test_array_interface.cu @@ -32,7 +32,7 @@ TEST(ArrayInterface, Stream) { dh::caching_device_vector out(1, 0); uint64_t dur = 1e9; dh::LaunchKernel{1, 1, 0, stream}(SleepForTest, out.data().get(), dur); - ArrayInterface arr(arr_str); + ArrayInterface<2> arr(arr_str); auto t = out[0]; CHECK_GE(t, dur); diff --git a/tests/cpp/data/test_iterative_device_dmatrix.cu b/tests/cpp/data/test_iterative_device_dmatrix.cu index cb64a3b5cdb2..27f6b0b3ffe9 100644 --- a/tests/cpp/data/test_iterative_device_dmatrix.cu +++ b/tests/cpp/data/test_iterative_device_dmatrix.cu @@ -103,7 +103,7 @@ TEST(IterativeDeviceDMatrix, RowMajor) { auto j_interface = Json::Load({interface_str.c_str(), interface_str.size()}); - ArrayInterface loaded {get(j_interface)}; + ArrayInterface<2> loaded {get(j_interface)}; std::vector h_data(cols * rows); common::Span s_data{static_cast(loaded.data), cols * rows}; dh::CopyDeviceSpanToVector(&h_data, s_data); @@ -128,7 +128,7 @@ TEST(IterativeDeviceDMatrix, RowMajorMissing) { std::string interface_str = iter.AsArray(); auto j_interface = Json::Load({interface_str.c_str(), interface_str.size()}); - ArrayInterface loaded {get(j_interface)}; + ArrayInterface<2> loaded {get(j_interface)}; std::vector h_data(cols * rows); common::Span s_data{static_cast(loaded.data), cols * rows}; dh::CopyDeviceSpanToVector(&h_data, s_data); diff --git a/tests/cpp/data/test_metainfo.cc b/tests/cpp/data/test_metainfo.cc index bb5452a56d28..c0249949eec1 100644 --- a/tests/cpp/data/test_metainfo.cc +++ b/tests/cpp/data/test_metainfo.cc @@ -122,7 +122,10 @@ TEST(MetaInfo, SaveLoadBinary) { EXPECT_EQ(inforead.labels_.HostVector(), info.labels_.HostVector()); EXPECT_EQ(inforead.group_ptr_, info.group_ptr_); EXPECT_EQ(inforead.weights_.HostVector(), info.weights_.HostVector()); - EXPECT_EQ(inforead.base_margin_.HostVector(), info.base_margin_.HostVector()); + + auto orig_margin = info.base_margin_.View(xgboost::GenericParameter::kCpuId); + auto read_margin = inforead.base_margin_.View(xgboost::GenericParameter::kCpuId); + EXPECT_TRUE(std::equal(orig_margin.cbegin(), orig_margin.cend(), read_margin.cbegin())); EXPECT_EQ(inforead.feature_type_names.size(), kCols); EXPECT_EQ(inforead.feature_types.Size(), kCols); @@ -257,7 +260,7 @@ TEST(MetaInfo, Validate) { auto arr_interface = xgboost::GetArrayInterface(&d_groups, 64, 1); std::string arr_interface_str; xgboost::Json::Dump(arr_interface, &arr_interface_str); - EXPECT_THROW(info.SetInfo("group", arr_interface_str), dmlc::Error); + EXPECT_THROW(info.SetInfo("group", xgboost::StringView{arr_interface_str}), dmlc::Error); #endif // defined(XGBOOST_USE_CUDA) } diff --git a/tests/cpp/data/test_metainfo.cu b/tests/cpp/data/test_metainfo.cu index 090374b913d6..9c278050ff3a 100644 --- a/tests/cpp/data/test_metainfo.cu +++ b/tests/cpp/data/test_metainfo.cu @@ -60,9 +60,9 @@ TEST(MetaInfo, FromInterface) { } info.SetInfo("base_margin", str.c_str()); - auto const& h_base_margin = info.base_margin_.HostVector(); + auto const h_base_margin = info.base_margin_.View(0); for (size_t i = 0; i < d_data.size(); ++i) { - ASSERT_EQ(h_base_margin[i], d_data[i]); + ASSERT_EQ(h_base_margin(0), d_data[i]); } thrust::device_vector d_group_data; diff --git a/tests/cpp/data/test_simple_dmatrix.cc b/tests/cpp/data/test_simple_dmatrix.cc index f777b00e244c..51b2763fb656 100644 --- a/tests/cpp/data/test_simple_dmatrix.cc +++ b/tests/cpp/data/test_simple_dmatrix.cc @@ -253,8 +253,8 @@ TEST(SimpleDMatrix, Slice) { std::iota(lower.begin(), lower.end(), 0.0f); std::iota(upper.begin(), upper.end(), 1.0f); - auto& margin = p_m->Info().base_margin_.HostVector(); - margin.resize(kRows * kClasses); + auto& margin = p_m->Info().base_margin_; + margin = linalg::Tensor{{kRows, kClasses}, GenericParameter::kCpuId}; std::array ridxs {1, 3, 5}; std::unique_ptr out { p_m->Slice(ridxs) }; @@ -284,10 +284,10 @@ TEST(SimpleDMatrix, Slice) { ASSERT_EQ(p_m->Info().weights_.HostVector().at(ridx), out->Info().weights_.HostVector().at(i)); - auto& out_margin = out->Info().base_margin_.HostVector(); + auto out_margin = out->Info().base_margin_.View(GenericParameter::kCpuId); + auto in_margin = margin.View(GenericParameter::kCpuId); for (size_t j = 0; j < kClasses; ++j) { - auto in_beg = ridx * kClasses; - ASSERT_EQ(out_margin.at(i * kClasses + j), margin.at(in_beg + j)); + ASSERT_EQ(out_margin(ridx, j), in_margin(ridx, j)); } } } diff --git a/tests/cpp/predictor/test_gpu_predictor.cu b/tests/cpp/predictor/test_gpu_predictor.cu index ad1083f9161b..b36df742da4f 100644 --- a/tests/cpp/predictor/test_gpu_predictor.cu +++ b/tests/cpp/predictor/test_gpu_predictor.cu @@ -108,7 +108,9 @@ TEST(GPUPredictor, ExternalMemoryTest) { dmats.push_back(CreateSparsePageDMatrix(8000)); for (const auto& dmat: dmats) { - dmat->Info().base_margin_.Resize(dmat->Info().num_row_ * n_classes, 0.5); + dmat->Info().base_margin_ = + linalg::Tensor{{dmat->Info().num_row_, static_cast(n_classes)}, 0}; + dmat->Info().base_margin_.Data()->Fill(0.5); PredictionCacheEntry out_predictions; gpu_predictor->InitOutPredictions(dmat->Info(), &out_predictions.predictions, model); gpu_predictor->PredictBatch(dmat.get(), &out_predictions, model, 0);