From 008f198b0d9d274f1752fdc10018f55e5eefb4d2 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 02:32:46 +0800 Subject: [PATCH 1/7] Extend array interface to handle ndarray. The `ArrayInterface` class is extended to support multi-dim array inputs. Previously this class handles only 2-dim (vector is also matrix). This PR specifies the expected dimension at compile-time and the array interface can perform various checks automatically for input data. Also, adapters like CSR are more rigorous about their input. Lastly, row vector and column vector are handled without intervention from the caller. --- include/xgboost/data.h | 8 +- .../xgboost4j-gpu/src/native/xgboost4j-gpu.cu | 30 +- src/data/adapter.h | 80 ++- src/data/array_interface.cu | 53 +- src/data/array_interface.h | 489 +++++++++++------- src/data/data.cc | 18 +- src/data/data.cu | 188 +++---- src/data/device_adapter.cuh | 47 +- src/data/file_iterator.h | 15 +- src/data/validation.h | 40 ++ tests/cpp/data/test_adapter.cc | 7 +- tests/cpp/data/test_array_interface.cc | 66 ++- tests/cpp/data/test_array_interface.cu | 15 +- tests/cpp/data/test_array_interface.h | 10 +- .../cpp/data/test_iterative_device_dmatrix.cu | 4 +- tests/cpp/data/test_metainfo.cu | 2 +- tests/cpp/helpers.h | 2 +- 17 files changed, 636 insertions(+), 438 deletions(-) create mode 100644 src/data/validation.h diff --git a/include/xgboost/data.h b/include/xgboost/data.h index 1678f4b1f4f1..a5210890da57 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 @@ -157,7 +159,7 @@ class MetaInfo { * * Right now only 1 column is permitted. */ - void SetInfo(const char* key, std::string const& interface_str); + void SetInfo(StringView key, std::string const& interface_str); void GetInfo(char const* key, bst_ulong* out_len, DataType dtype, const void** out_dptr) const; diff --git a/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu index de7a1fc41495..4ecf8b0f1da1 100644 --- a/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu +++ b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu @@ -35,13 +35,12 @@ template T CheckJvmCall(T const &v, JNIEnv *jenv) { } template -void CopyColumnMask(xgboost::ArrayInterface const &interface, +void CopyColumnMask(xgboost::ArrayInterface<1> const &interface, std::vector const &columns, cudaMemcpyKind kind, size_t c, VCont *p_mask, Json *p_out, cudaStream_t stream) { auto &mask = *p_mask; auto &out = *p_out; - auto size = sizeof(typename VCont::value_type) * interface.num_rows * - interface.num_cols; + auto size = sizeof(typename VCont::value_type) * interface.n; mask.resize(size); CHECK(RawPtr(mask)); CHECK(size); @@ -67,11 +66,11 @@ void CopyColumnMask(xgboost::ArrayInterface const &interface, LOG(FATAL) << "Invalid shape of mask"; } out["mask"]["typestr"] = String(" -void CopyInterface(std::vector &interface_arr, +void CopyInterface(std::vector> &interface_arr, std::vector const &columns, cudaMemcpyKind kind, std::vector *p_data, std::vector *p_mask, std::vector *p_out, cudaStream_t stream) { @@ -81,7 +80,7 @@ void CopyInterface(std::vector &interface_arr, for (size_t c = 0; c < interface_arr.size(); ++c) { auto &interface = interface_arr.at(c); size_t element_size = interface.ElementSize(); - size_t size = element_size * interface.num_rows * interface.num_cols; + size_t size = element_size * interface.n; auto &data = (*p_data)[c]; auto &mask = (*p_mask)[c]; @@ -95,14 +94,13 @@ void CopyInterface(std::vector &interface_arr, Json{Boolean{false}}}; out["data"] = Array(std::move(j_data)); - out["shape"] = Array(std::vector{Json(Integer(interface.num_rows)), - Json(Integer(interface.num_cols))}); + out["shape"] = Array(std::vector{Json(Integer(interface.Shape(0)))}); if (interface.valid.Data()) { CopyColumnMask(interface, columns, kind, c, &mask, &out, stream); } out["typestr"] = String(" *out, cudaStream_t auto &j_interface = *p_interface; CHECK_EQ(get(j_interface).size(), 1); auto object = get(get(j_interface)[0]); - ArrayInterface interface(object); - out->resize(interface.num_rows); + ArrayInterface<1> interface(object); + out->resize(interface.Shape(0)); size_t element_size = interface.ElementSize(); - size_t size = element_size * interface.num_rows; + size_t size = element_size * interface.n; dh::safe_cuda(cudaMemcpyAsync(RawPtr(*out), interface.data, size, cudaMemcpyDeviceToDevice, stream)); j_interface[0]["data"][0] = reinterpret_cast(RawPtr(*out)); @@ -285,11 +283,11 @@ class DataIteratorProxy { Json features = json_interface["features_str"]; auto json_columns = get(features); - std::vector interfaces; + std::vector> interfaces; // Stage the data for (auto &json_col : json_columns) { - auto column = ArrayInterface(get(json_col)); + auto column = ArrayInterface<1>(get(json_col)); interfaces.emplace_back(column); } Json::Dump(features, &interface_str); @@ -342,9 +340,9 @@ class DataIteratorProxy { // Data auto const &json_interface = host_columns_.at(it_)->interfaces; - std::vector in; + std::vector> in; for (auto interface : json_interface) { - auto column = ArrayInterface(get(interface)); + auto column = ArrayInterface<1>(get(interface)); in.emplace_back(column); } std::vector out; diff --git a/src/data/adapter.h b/src/data/adapter.h index 27da8c6e3b36..4214171e8b7e 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,42 @@ 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_, TypedIndex{indices_}(offset_ + idx), values_(offset_ + idx)}; } size_t Size() const { - return values_.num_rows * values_.num_cols; + return values_.Shape(0); } }; @@ -339,17 +338,16 @@ 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} { } size_t NumRows() const { - size_t size = indptr_.num_rows * indptr_.num_cols; + size_t size = indptr_.Shape(0); size = size == 0 ? 0 : size - 1; return size; } @@ -357,19 +355,19 @@ 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_no_stride = TypedIndex{indptr_}(idx); + auto end_no_stride = TypedIndex{indptr_}(idx + 1); auto indices = indices_; auto values = values_; + // Slice indices and values, stride remains unchanged since this is slicing by + // specific index. + auto offset = indices.strides[0] * begin_no_stride; - values.num_cols = end_offset - begin_offset; - values.num_rows = 1; + indices.shape[0] = end_no_stride - begin_no_stride; + values.shape[0] = end_no_stride - begin_no_stride; - indices.num_cols = values.num_cols; - indices.num_rows = values.num_rows; - - return Line{indices, values, idx, begin_offset}; + return Line{indices, values, idx, offset}; } }; @@ -391,7 +389,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(0); size = size == 0 ? 0 : size - 1; return size; } @@ -399,9 +397,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.cu b/src/data/array_interface.cu index def4de195523..aad12e6f9dc8 100644 --- a/src/data/array_interface.cu +++ b/src/data/array_interface.cu @@ -7,15 +7,50 @@ namespace xgboost { void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { switch (stream) { - case 0: - LOG(FATAL) << "Invalid stream ID in array interface: " << stream; - case 1: - // default legacy stream - break; - case 2: - // default per-thread stream - default: - dh::safe_cuda(cudaStreamSynchronize(reinterpret_cast(stream))); + case 0: + /** + * disallowed by the `__cuda_array_interface__`. Quote: + * + * This is disallowed as it would be ambiguous between None and the default + * stream, and also between the legacy and per-thread default streams. Any use + * case where 0 might be given should either use None, 1, or 2 instead for + * clarity. + */ + LOG(FATAL) << "Invalid stream ID in array interface: " << stream; + case 1: + // default legacy stream + break; + case 2: + // default per-thread stream + default: + dh::safe_cuda(cudaStreamSynchronize(reinterpret_cast(stream))); + } +} + +bool ArrayInterfaceHandler::IsCudaPtr(void const* ptr) { + if (!ptr) { + return false; + } + cudaPointerAttributes attr; + auto err = cudaPointerGetAttributes(&attr, ptr); + // reset error + CHECK_EQ(err, cudaGetLastError()); + if (err == cudaErrorInvalidValue) { + // CUDA < 11 + return false; + } else if (err == cudaSuccess) { + // CUDA >= 11 + switch (attr.type) { + case cudaMemoryTypeUnregistered: + case cudaMemoryTypeHost: + return false; + default: + return true; + } + return true; + } else { + // other errors, `cudaErrorNoDevice`, `cudaErrorInsufficientDriver` etc. + return false; } } } // namespace xgboost diff --git a/src/data/array_interface.h b/src/data/array_interface.h index 6524f4512407..a83142c4d5fa 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 *Contiguous() { return "Memory should be contiguous."; } + 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() { - return "Only version <= 3 of " - "`__cuda_array_interface__/__array_interface__' are supported."; + static char const *Version() { + return "Only version <= 3 of `__cuda_array_interface__' and `__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,38 +95,25 @@ struct ArrayInterfaceErrors { // object and turn it into an array (for cupy and numba). class ArrayInterfaceHandler { public: - template - static constexpr char TypeChar() { - return - (std::is_floating_point::value ? 'f' : - (std::is_integral::value ? - (std::is_signed::value ? 'i' : 'u') : '\0')); - } + enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 }; 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"; } - auto stream_it = array.find("stream"); - if (stream_it != array.cend() && !IsA(stream_it->second)) { - // is cuda, check the version. - if (get(version_it->second) > 3) { - LOG(FATAL) << ArrayInterfaceErrors::Version(); - } + if (get(version_it->second) > 3) { + LOG(FATAL) << ArrayInterfaceErrors::Version(); } if (array.find("typestr") == array.cend()) { @@ -149,12 +135,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,8 +172,8 @@ 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(get(strides.at(0)), type_length) << ArrayInterfaceErrors::Contigious(); + CHECK_EQ(strides.size(), 1) << ArrayInterfaceErrors::Dimension(1); + CHECK_EQ(get(strides.at(0)), type_length) << ArrayInterfaceErrors::Contiguous(); } s_mask = {p_mask, span_size}; @@ -195,77 +181,213 @@ 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)))}; + /** + * \brief Handle vector inputs. For higher dimension, we require strictly correct shape. + */ + template + static void HandleRowVector(std::vector const &shape, std::vector *p_out) { + auto &out = *p_out; + if (shape.size() == 2 && D == 1) { + auto m = shape[0]; + auto n = shape[1]; + CHECK(m == 1 || n == 1); + if (m == 1) { + // keep the number of columns + out[0] = out[1]; + out.resize(1); + } else if (n == 1) { + // keep the number of rows. + out.resize(1); + } + // when both m and n are 1, above logic keeps the column. + // when neither m nor n is 1, caller should throw an error about Dimension. } } - 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 void ExtractShape(std::map const &array, size_t (&out_shape)[D]) { + auto const &j_shape = get(array.at("shape")); + std::vector shape_arr(j_shape.size(), 0); + std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(), + [](Json in) { return get(in); }); + // handle column vector vs. row vector + HandleRowVector(shape_arr, &shape_arr); + // Copy shape. + size_t i; + for (i = 0; i < shape_arr.size(); ++i) { + CHECK_LT(i, D) << ArrayInterfaceErrors::Dimension(D); + out_shape[i] = shape_arr[i]; } + // Fill the remaining dimensions + std::fill(out_shape + i, out_shape + D, 1); + } - auto valid = rows * (*stride_r) + cols * (*stride_c) >= (rows * cols); - CHECK(valid) << "Invalid strides in array." - << " strides: (" << (*stride_r) << "," << (*stride_c) - << "), shape: (" << rows << ", " << cols << ")"; + /** + * \brief Extracts the optiona `strides' field and returns whether the array is c-contiguous. + */ + 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"); + // No stride is provided + if (strides_it == array.cend() || IsA(strides_it->second)) { + // No stride is provided, we can calculate it from shape. + linalg::detail::CalcStride(shape, stride); + // Quote: + // + // strides: Either None to indicate a C-style contiguous array or a Tuple of + // strides which provides the number of bytes + return true; + } + // Get shape, we need to make changes to handle row vector, so some duplicated code + // from `ExtractShape` for copying out the shape. + auto const &j_shape = get(array.at("shape")); + std::vector shape_arr(j_shape.size(), 0); + std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(), + [](Json in) { return get(in); }); + // Get stride + auto const &j_strides = get(strides_it->second); + CHECK_EQ(j_strides.size(), j_shape.size()) << "stride and shape don't match."; + std::vector stride_arr(j_strides.size(), 0); + std::transform(j_strides.cbegin(), j_strides.cend(), stride_arr.begin(), + [](Json in) { return get(in); }); + + // Handle column vector vs. row vector + HandleRowVector(shape_arr, &stride_arr); + size_t i; + for (i = 0; i < stride_arr.size(); ++i) { + // If one of the dim has shape 0 then total size is 0, stride is meaningless, but we + // set it to 0 here just to be consistent + CHECK_LT(i, D) << ArrayInterfaceErrors::Dimension(D); + // We use number of items instead of number of bytes + stride[i] = stride_arr[i] / itemsize; + } + std::fill(stride + i, stride + D, 1); + // If the stride can be calculated from shape then it's contiguous. + size_t stride_tmp[D]; + linalg::detail::CalcStride(shape, stride_tmp); + return std::equal(stride_tmp, stride_tmp + D, stride); } - 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; } - + /** + * \brief Whether the ptr is allocated by CUDA. + */ + static bool IsCudaPtr(void const *ptr); + /** + * \brief Sync the CUDA stream. + */ static void SyncCudaStream(int64_t stream); }; +/** + * Dispatch compile time type to runtime type. + */ +template +struct ToDType; +// float +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF8; +}; +template +struct ToDType::value && sizeof(long double) == 16>> { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF16; +}; +// uint +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU1; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU2; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU4; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU8; +}; +// int +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI1; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI2; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI4; +}; +template <> +struct ToDType { + static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI8; +}; + #if !defined(XGBOOST_USE_CUDA) -inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { - common::AssertGPUSupport(); -} +inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { common::AssertGPUSupport(); } +inline bool ArrayInterfaceHandler::IsCudaPtr(void const* ptr) { return false; } #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_contiguous = 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); + static_assert(allow_mask ? D == 1 : D >= 1, "Masked ndarray is not supported."); if (allow_mask) { common::Span s_mask; size_t n_bits = ArrayInterfaceHandler::ExtractMask(array, &s_mask); @@ -273,18 +395,13 @@ class ArrayInterface { 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 +409,149 @@ 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) { - auto jinterface = Json::Load(str); - if (IsA(jinterface)) { - this->Initialize(get(jinterface), allow_mask); + explicit ArrayInterface(Json const &array) { + if (IsA(array)) { + this->Initialize(get(array)); return; } - if (IsA(jinterface)) { - CHECK_EQ(get(jinterface).size(), 1) + if (IsA(array)) { + CHECK_EQ(get(array).size(), 1) << "Column: " << ArrayInterfaceErrors::Dimension(1); - this->Initialize(get(get(jinterface)[0]), allow_mask); + this->Initialize(get(get(array)[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; + explicit ArrayInterface(std::string const &str) : ArrayInterface{StringView{str}} {} - stride_row = std::max(stride_row, stride_col); - stride_col = 1; - } + explicit ArrayInterface(StringView str) : ArrayInterface{Json::Load(str)} {} 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 { + static_assert(sizeof...(index) <= D, "Invalid index."); + return this->DispatchCall([=](auto const *p_values) -> T { + size_t offset = linalg::detail::Offset<0ul>(strides, 0ul, index...); + return static_cast(p_values[offset]); + }); } + // Used only by columnar format. 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; + // Array stride + size_t strides[D]{0}; + // Array shape + size_t shape[D]{0}; + // Type earsed pointer referencing the data. + void *data; + // Total number of items + size_t n; + // Whether the memory is c-contiguous + bool is_contiguous {false}; + // RTTI + 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('>'); +/** + * \brief Helper for type casting. + */ +template +struct TypedIndex { + ArrayInterface const &array; + template + XGBOOST_DEVICE T operator()(I &&...ind) const { + static_assert(sizeof...(ind) <= D, "Invalid index."); + return array.template operator()(ind...); } - 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; +}; + +template +inline void CheckArrayInterface(StringView key, ArrayInterface const &array) { + CHECK(!array.valid.Data()) << "Meta info " << key << " should be dense, found validity mask"; } } // namespace xgboost #endif // XGBOOST_DATA_ARRAY_INTERFACE_H_ diff --git a/src/data/data.cc b/src/data/data.cc index acf6e47b9d21..a6b76ee2a39c 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 @@ -24,6 +24,7 @@ #include "../data/iterative_device_dmatrix.h" #include "file_iterator.h" +#include "validation.h" #include "./sparse_page_source.h" #include "./sparse_page_dmatrix.h" @@ -337,17 +338,6 @@ inline bool MetaTryLoadFloatInfo(const std::string& fname, return true; } -void ValidateQueryGroup(std::vector const &group_ptr_) { - bool valid_query_group = true; - for (size_t i = 1; i < group_ptr_.size(); ++i) { - valid_query_group = valid_query_group && group_ptr_[i] >= group_ptr_[i - 1]; - if (!valid_query_group) { - break; - } - } - CHECK(valid_query_group) << "Invalid group structure."; -} - // macro to dispatch according to specified pointer types #define DISPATCH_CONST_PTR(dtype, old_ptr, cast_ptr, proc) \ switch (dtype) { \ @@ -398,7 +388,7 @@ void MetaInfo::SetInfo(const char* key, const void* dptr, DataType dtype, size_t for (size_t i = 1; i < group_ptr_.size(); ++i) { group_ptr_[i] = group_ptr_[i - 1] + group_ptr_[i]; } - ValidateQueryGroup(group_ptr_); + data::ValidateQueryGroup(group_ptr_); } else if (!std::strcmp(key, "qid")) { std::vector query_ids(num, 0); DISPATCH_CONST_PTR(dtype, dptr, cast_dptr, @@ -632,7 +622,7 @@ void MetaInfo::Validate(int32_t device) const { } #if !defined(XGBOOST_USE_CUDA) -void MetaInfo::SetInfo(const char * c_key, std::string const& interface_str) { +void MetaInfo::SetInfo(StringView key, std::string const& interface_str) { common::AssertGPUSupport(); } #endif // !defined(XGBOOST_USE_CUDA) diff --git a/src/data/data.cu b/src/data/data.cu index aee62d1b7fad..d2807362896a 100644 --- a/src/data/data.cu +++ b/src/data/data.cu @@ -9,84 +9,81 @@ #include "xgboost/json.h" #include "array_interface.h" #include "../common/device_helpers.cuh" +#include "../common/linalg_op.cuh" #include "device_adapter.cuh" #include "simple_dmatrix.h" +#include "validation.h" namespace xgboost { - -void CopyInfoImpl(ArrayInterface column, HostDeviceVector* out) { - auto SetDeviceToPtr = [](void* ptr) { - cudaPointerAttributes attr; - dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); - int32_t ptr_device = attr.device; - if (ptr_device >= 0) { - dh::safe_cuda(cudaSetDevice(ptr_device)); - } - return ptr_device; - }; - auto ptr_device = SetDeviceToPtr(column.data); - - if (column.num_rows == 0) { - return; - } - out->SetDevice(ptr_device); - - size_t size = column.num_rows * column.num_cols; - 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) { +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; } -} // anonymous namespace -void CopyGroupInfoImpl(ArrayInterface column, std::vector* out) { - CHECK(column.type != ArrayInterface::kF4 && column.type != ArrayInterface::kF8) +template +void CopyTensorInfoImpl(Json arr_interface, linalg::Tensor* p_out) { + ArrayInterface array(arr_interface); + if (array.n == 0) { + return; + } + CHECK(array.valid.Size() == 0) << "Meta info like label or weight can not have missing value."; + auto ptr_device = SetDeviceToPtr(array.data); + + if (array.is_contiguous && array.type == ToDType::kType) { + p_out->ModifyInplace([&](HostDeviceVector* data, common::Span shape) { + // set shape + std::copy(array.shape, array.shape + D, shape.data()); + // set data + data->SetDevice(ptr_device); + data->Resize(array.n); + dh::safe_cuda(cudaMemcpyAsync(data->DevicePointer(), array.data, array.n * sizeof(T), + cudaMemcpyDefault)); + }); + return; + } + p_out->SetDevice(ptr_device); + p_out->Reshape(array.shape); + auto t = p_out->View(ptr_device); + linalg::ElementWiseKernelDevice(t, [=] __device__(size_t i, T) { + return linalg::detail::Apply(TypedIndex{array}, linalg::UnravelIndex(i, array.shape)); + }); +} + +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); - auto d_tmp = temp.data(); + dh::TemporaryArray temp(column.Shape(0)); + auto d_tmp = temp.data().get(); - 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] = TypedIndex{column}(idx); }); + 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); + thrust::make_counting_iterator(0ul), [array_interface] __device__(size_t i) { + return TypedIndex{array_interface}(i); }); 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) { + auto typed = TypedIndex{array_interface}; + if (typed(i) > typed(i + 1)) { d_flag[0] = false; } }); @@ -95,16 +92,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(); @@ -115,77 +112,56 @@ void CopyQidImpl(ArrayInterface array_interface, thrust::copy(cnt.begin(), cnt.begin() + h_num_runs_out, group_ptr_.begin() + 1); } +} // namespace -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 - -void ValidateQueryGroup(std::vector const &group_ptr_); - -void MetaInfo::SetInfo(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}; - - CHECK(!array_interface.valid.Data()) - << "Meta info " << key << " should be dense, found validity mask"; - if (array_interface.num_rows == 0) { - return; - } - +void MetaInfo::SetInfo(StringView key, std::string const& interface_str) { + Json array = Json::Load(StringView{interface_str}); + // multi-dim float info if (key == "base_margin") { - CopyInfoImpl(array_interface, &base_margin_); + // FIXME(jiamingy): This is temporary until #7405 can be fully merged + linalg::Tensor t; + CopyTensorInfoImpl(array, &t); + base_margin_ = std::move(*t.Data()); return; } - - 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(); + // uint info + if (key == "group") { + auto array_interface{ArrayInterface<1>(array)}; + CopyGroupInfoImpl(array_interface, &group_ptr_); + data::ValidateQueryGroup(group_ptr_); + return; + } else if (key == "qid") { + auto array_interface{ArrayInterface<1>(array)}; + CopyQidImpl(array_interface, &group_ptr_); + data::ValidateQueryGroup(group_ptr_); + return; } + // float info + linalg::Tensor t; + CopyTensorInfoImpl(array, &t); if (key == "label") { - CopyInfoImpl(array_interface, &labels_); + this->labels_ = std::move(*t.Data()); 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(), data::LabelsCheck{}); CHECK(valid) << "Label contains NaN, infinity or a value too large."; } else if (key == "weight") { - CopyInfoImpl(array_interface, &weights_); + this->weights_ = std::move(*t.Data()); 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(), data::WeightsCheck{}); CHECK(valid) << "Weights must be positive values."; - } else if (key == "group") { - CopyGroupInfoImpl(array_interface, &group_ptr_); - ValidateQueryGroup(group_ptr_); - return; - } else if (key == "qid") { - CopyQidImpl(array_interface, &group_ptr_); - return; } else if (key == "label_lower_bound") { - CopyInfoImpl(array_interface, &labels_lower_bound_); - return; + this->labels_lower_bound_ = std::move(*t.Data()); } else if (key == "label_upper_bound") { - CopyInfoImpl(array_interface, &labels_upper_bound_); - return; + this->labels_upper_bound_ = std::move(*t.Data()); } else if (key == "feature_weights") { - CopyInfoImpl(array_interface, &feature_weights); + this->feature_weights = std::move(*t.Data()); auto d_feature_weights = feature_weights.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(), data::WeightsCheck{}); CHECK(valid) << "Feature weight must be greater than 0."; - return; } else { - LOG(FATAL) << "Unknown metainfo: " << key; + LOG(FATAL) << "Unknown key for MetaInfo: " << key; } } diff --git a/src/data/device_adapter.cuh b/src/data/device_adapter.cuh index 628878f319f1..d1bda280a7d5 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) : 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,12 @@ 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)); + 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 +134,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 +142,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 +166,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 +176,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..70a5d51c30b9 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).ArrayInterfaceStr(); + values_ = MakeVec(row_block_.value, row_block_.offset[row_block_.size]).ArrayInterfaceStr(); + indices_ = MakeVec(row_block_.index, row_block_.offset[row_block_.size]).ArrayInterfaceStr(); - 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/validation.h b/src/data/validation.h new file mode 100644 index 000000000000..6d3701114886 --- /dev/null +++ b/src/data/validation.h @@ -0,0 +1,40 @@ +/*! + * Copyright 2021 by XGBoost Contributors + */ +#ifndef XGBOOST_DATA_VALIDATION_H_ +#define XGBOOST_DATA_VALIDATION_H_ +#include +#include + +#include "xgboost/base.h" +#include "xgboost/logging.h" + +namespace xgboost { +namespace data { +struct LabelsCheck { + XGBOOST_DEVICE bool operator()(float y) { +#if defined(__CUDA_ARCH__) + return ::isnan(y) || ::isinf(y); +#else + return std::isnan(y) || std::isinf(y); +#endif + } +}; + +struct WeightsCheck { + XGBOOST_DEVICE bool operator()(float w) { return LabelsCheck{}(w) || w < 0; } // NOLINT +}; + +inline void ValidateQueryGroup(std::vector const &group_ptr_) { + bool valid_query_group = true; + for (size_t i = 1; i < group_ptr_.size(); ++i) { + valid_query_group = valid_query_group && group_ptr_[i] >= group_ptr_[i - 1]; + if (XGBOOST_EXPECT(!valid_query_group, false)) { + break; + } + } + CHECK(valid_query_group) << "Invalid group structure."; +} +} // namespace data +} // namespace xgboost +#endif // XGBOOST_DATA_VALIDATION_H_ diff --git a/tests/cpp/data/test_adapter.cc b/tests/cpp/data/test_adapter.cc index ccb19de71a74..c6de226194f6 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()).ArrayInterfaceStr(); + auto values_arr = MakeVec(values.HostPointer(), values.Size()).ArrayInterfaceStr(); + auto indices_arr = MakeVec(indices.HostPointer(), indices.Size()).ArrayInterfaceStr(); 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..8efd30eb75b2 100644 --- a/tests/cpp/data/test_array_interface.cc +++ b/tests/cpp/data/test_array_interface.cc @@ -11,21 +11,22 @@ 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::string u64_arr_str{linalg::TensorView{ + u64_storage.ConstHostSpan(), {kRows, kCols}, GenericParameter::kCpuId} + .ArrayInterfaceStr()}; 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 +39,22 @@ TEST(ArrayInterface, Error) { Json(Boolean(false))}; auto const& column_obj = get(column); - std::pair shape{kRows, kCols}; std::string typestr{"(1)); + EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n), dmlc::Error); + column["version"] = 3; // 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 +63,52 @@ 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); + } +} + +TEST(ArrayInterface, ToDType) { + static_assert(ToDType::kType == ArrayInterfaceHandler::kF4, ""); + static_assert(ToDType::kType == ArrayInterfaceHandler::kF8, ""); + + static_assert(ToDType::kType == ArrayInterfaceHandler::kU4, ""); + static_assert(ToDType::kType == ArrayInterfaceHandler::kU8, ""); + + static_assert(ToDType::kType == ArrayInterfaceHandler::kI4, ""); + static_assert(ToDType::kType == ArrayInterfaceHandler::kI8, ""); +} } // namespace xgboost diff --git a/tests/cpp/data/test_array_interface.cu b/tests/cpp/data/test_array_interface.cu index 75923e74ba1a..c8e07852534b 100644 --- a/tests/cpp/data/test_array_interface.cu +++ b/tests/cpp/data/test_array_interface.cu @@ -32,11 +32,24 @@ 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); cudaStreamDestroy(stream); } + +TEST(ArrayInterface, Ptr) { + std::vector h_data(10); + ASSERT_FALSE(ArrayInterfaceHandler::IsCudaPtr(h_data.data())); + dh::safe_cuda(cudaGetLastError()); + + dh::device_vector d_data(10); + ASSERT_TRUE(ArrayInterfaceHandler::IsCudaPtr(d_data.data().get())); + dh::safe_cuda(cudaGetLastError()); + + ASSERT_FALSE(ArrayInterfaceHandler::IsCudaPtr(nullptr)); + dh::safe_cuda(cudaGetLastError()); +} } // namespace xgboost diff --git a/tests/cpp/data/test_array_interface.h b/tests/cpp/data/test_array_interface.h index 7872a9507aa5..78bce76f53e7 100644 --- a/tests/cpp/data/test_array_interface.h +++ b/tests/cpp/data/test_array_interface.h @@ -19,6 +19,7 @@ Json GenerateDenseColumn(std::string const& typestr, size_t kRows, std::vector j_shape {Json(Integer(static_cast(kRows)))}; column["shape"] = Array(j_shape); column["strides"] = Array(std::vector{Json(Integer(static_cast(sizeof(T))))}); + column["stream"] = nullptr; d_data.resize(kRows); thrust::sequence(thrust::device, d_data.begin(), d_data.end(), 0.0f, 2.0f); @@ -30,7 +31,7 @@ Json GenerateDenseColumn(std::string const& typestr, size_t kRows, Json(Boolean(false))}; column["data"] = j_data; - column["version"] = Integer(static_cast(1)); + column["version"] = 3; column["typestr"] = String(typestr); return column; } @@ -43,6 +44,7 @@ Json GenerateSparseColumn(std::string const& typestr, size_t kRows, std::vector j_shape {Json(Integer(static_cast(kRows)))}; column["shape"] = Array(j_shape); column["strides"] = Array(std::vector{Json(Integer(static_cast(sizeof(T))))}); + column["stream"] = nullptr; d_data.resize(kRows); for (size_t i = 0; i < d_data.size(); ++i) { @@ -56,7 +58,7 @@ Json GenerateSparseColumn(std::string const& typestr, size_t kRows, Json(Boolean(false))}; column["data"] = j_data; - column["version"] = Integer(static_cast(1)); + column["version"] = 3; column["typestr"] = String(typestr); return column; } @@ -75,9 +77,9 @@ Json Generate2dArrayInterface(int rows, int cols, std::string typestr, Json(Integer(reinterpret_cast(data.data().get()))), Json(Boolean(false))}; array_interface["data"] = j_data; - array_interface["version"] = Integer(static_cast(1)); + array_interface["version"] = 3; array_interface["typestr"] = String(typestr); + array_interface["stream"] = nullptr; return array_interface; } - } // namespace xgboost 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.cu b/tests/cpp/data/test_metainfo.cu index 090374b913d6..205844a5e961 100644 --- a/tests/cpp/data/test_metainfo.cu +++ b/tests/cpp/data/test_metainfo.cu @@ -1,4 +1,4 @@ -/*! Copyright 2019 by Contributors */ +/*! Copyright 2019-2021 by XGBoost Contributors */ #include #include diff --git a/tests/cpp/helpers.h b/tests/cpp/helpers.h index c424d65ced05..2ccbd8a77e83 100644 --- a/tests/cpp/helpers.h +++ b/tests/cpp/helpers.h @@ -198,7 +198,7 @@ Json GetArrayInterface(HostDeviceVector *storage, size_t rows, size_t cols) { array_interface["shape"][0] = rows; array_interface["shape"][1] = cols; - char t = ArrayInterfaceHandler::TypeChar(); + char t = linalg::detail::ArrayInterfaceHandler::TypeChar(); array_interface["typestr"] = String(std::string{"<"} + t + std::to_string(sizeof(T))); array_interface["version"] = 1; return array_interface; From ffbb9e3554f39f64f91d6a5dfed0d57816c4de9f Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 03:52:56 +0800 Subject: [PATCH 2/7] Polish. --- src/data/array_interface.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/data/array_interface.h b/src/data/array_interface.h index a83142c4d5fa..741ccf1f595f 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -116,11 +116,12 @@ class ArrayInterfaceHandler { LOG(FATAL) << ArrayInterfaceErrors::Version(); } - if (array.find("typestr") == array.cend()) { + auto typestr_it = array.find("typestr"); + if (typestr_it == array.cend()) { LOG(FATAL) << "Missing `typestr' field for array interface"; } - auto typestr = get(array.at("typestr")); + auto typestr = get(typestr_it->second); CHECK(typestr.size() == 3 || typestr.size() == 4) << ArrayInterfaceErrors::TypestrFormat(); if (array.find("shape") == array.cend()) { From aa6fd93627fc1407c677c1f9c87f2d5765b30f85 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 03:54:08 +0800 Subject: [PATCH 3/7] Redundant comment. --- src/data/array_interface.h | 1 - 1 file changed, 1 deletion(-) diff --git a/src/data/array_interface.h b/src/data/array_interface.h index 741ccf1f595f..5fd266844742 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -347,7 +347,6 @@ inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { common::Asse inline bool ArrayInterfaceHandler::IsCudaPtr(void const* ptr) { return false; } #endif // !defined(XGBOOST_USE_CUDA) -// A view over __array_interface__ /** * \brief A type erased view over __array_interface__ protocol defined by numpy * From b600597bf8efa51e1d517ead4b9805d0d7764446 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 04:02:07 +0800 Subject: [PATCH 4/7] Make const. --- src/common/device_helpers.cuh | 2 +- src/data/array_interface.h | 38 ++++++++++--------- src/data/data.cu | 2 +- .../cpp/data/test_iterative_device_dmatrix.cu | 4 +- 4 files changed, 24 insertions(+), 22 deletions(-) diff --git a/src/common/device_helpers.cuh b/src/common/device_helpers.cuh index f35af3b5e647..b316453f5d2c 100644 --- a/src/common/device_helpers.cuh +++ b/src/common/device_helpers.cuh @@ -136,7 +136,7 @@ inline ncclResult_t ThrowOnNcclError(ncclResult_t code, const char *file, } #endif -inline int32_t CudaGetPointerDevice(void* ptr) { +inline int32_t CudaGetPointerDevice(void const *ptr) { int32_t device = -1; cudaPointerAttributes attr; dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); diff --git a/src/data/array_interface.h b/src/data/array_interface.h index 5fd266844742..9dea2f4e44ea 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -91,19 +91,21 @@ struct ArrayInterfaceErrors { } }; -// TODO(trivialfis): Abstract this into a class that accept a json -// object and turn it into an array (for cupy and numba). +/** + * Utilities for consuming array interface. + */ class ArrayInterfaceHandler { public: enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 }; template static PtrType GetPtrFromArrayData(std::map const &obj) { - if (obj.find("data") == obj.cend()) { + auto data_it = obj.find("data"); + if (data_it == obj.cend()) { LOG(FATAL) << "Empty data passed in."; } auto p_data = reinterpret_cast( - static_cast(get(get(obj.at("data")).at(0)))); + static_cast(get(get(data_it->second).at(0)))); return p_data; } @@ -472,38 +474,38 @@ class ArrayInterface { using T = ArrayInterfaceHandler::Type; switch (type) { case T::kF4: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kF8: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); #ifdef __CUDA_ARCH__ case T::kF16: { // CUDA device code doesn't support long double. SPAN_CHECK(false); - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); } #else case T::kF16: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); #endif case T::kI1: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kI2: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kI4: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kI8: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kU1: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kU2: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kU4: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); case T::kU8: - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); } SPAN_CHECK(false); - return func(reinterpret_cast(data)); + return func(reinterpret_cast(data)); } XGBOOST_DEVICE size_t constexpr ElementSize() { @@ -527,7 +529,7 @@ class ArrayInterface { // Array shape size_t shape[D]{0}; // Type earsed pointer referencing the data. - void *data; + void const *data; // Total number of items size_t n; // Whether the memory is c-contiguous diff --git a/src/data/data.cu b/src/data/data.cu index d2807362896a..a9e333f31db1 100644 --- a/src/data/data.cu +++ b/src/data/data.cu @@ -16,7 +16,7 @@ namespace xgboost { namespace { -auto SetDeviceToPtr(void* ptr) { +auto SetDeviceToPtr(void const* ptr) { cudaPointerAttributes attr; dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr)); int32_t ptr_device = attr.device; diff --git a/tests/cpp/data/test_iterative_device_dmatrix.cu b/tests/cpp/data/test_iterative_device_dmatrix.cu index 27f6b0b3ffe9..0fc992f24187 100644 --- a/tests/cpp/data/test_iterative_device_dmatrix.cu +++ b/tests/cpp/data/test_iterative_device_dmatrix.cu @@ -105,7 +105,7 @@ TEST(IterativeDeviceDMatrix, RowMajor) { Json::Load({interface_str.c_str(), interface_str.size()}); ArrayInterface<2> loaded {get(j_interface)}; std::vector h_data(cols * rows); - common::Span s_data{static_cast(loaded.data), cols * rows}; + common::Span s_data{static_cast(loaded.data), cols * rows}; dh::CopyDeviceSpanToVector(&h_data, s_data); for(auto i = 0ull; i < rows * cols; i++) { @@ -130,7 +130,7 @@ TEST(IterativeDeviceDMatrix, RowMajorMissing) { Json::Load({interface_str.c_str(), interface_str.size()}); ArrayInterface<2> loaded {get(j_interface)}; std::vector h_data(cols * rows); - common::Span s_data{static_cast(loaded.data), cols * rows}; + common::Span s_data{static_cast(loaded.data), cols * rows}; dh::CopyDeviceSpanToVector(&h_data, s_data); h_data[1] = kMissing; h_data[5] = kMissing; From c08c547fabc15b1f81080bf6bd6ed4f135bdebf3 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 04:12:21 +0800 Subject: [PATCH 5/7] Initialization. --- src/data/array_interface.h | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/data/array_interface.h b/src/data/array_interface.h index 9dea2f4e44ea..f69b1dc22885 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -529,13 +529,13 @@ class ArrayInterface { // Array shape size_t shape[D]{0}; // Type earsed pointer referencing the data. - void const *data; + void const *data{nullptr}; // Total number of items - size_t n; + size_t n{0}; // Whether the memory is c-contiguous bool is_contiguous {false}; - // RTTI - ArrayInterfaceHandler::Type type; + // RTTI, initialized to the f16 to avoid masking potential bugs in initialization. + ArrayInterfaceHandler::Type type{ArrayInterfaceHandler::kF16}; }; /** From 66a0815b68825710d5b0488a2ff000267bec4a89 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 04:14:03 +0800 Subject: [PATCH 6/7] Format. --- src/data/array_interface.cu | 2 +- src/data/array_interface.h | 8 +++----- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/src/data/array_interface.cu b/src/data/array_interface.cu index aad12e6f9dc8..b1a80251ecc4 100644 --- a/src/data/array_interface.cu +++ b/src/data/array_interface.cu @@ -1,8 +1,8 @@ /*! * Copyright 2021 by Contributors */ -#include "array_interface.h" #include "../common/common.h" +#include "array_interface.h" namespace xgboost { void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { diff --git a/src/data/array_interface.h b/src/data/array_interface.h index f69b1dc22885..1b2545a57d62 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -346,7 +346,7 @@ struct ToDType { #if !defined(XGBOOST_USE_CUDA) inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { common::AssertGPUSupport(); } -inline bool ArrayInterfaceHandler::IsCudaPtr(void const* ptr) { return false; } +inline bool ArrayInterfaceHandler::IsCudaPtr(void const *ptr) { return false; } #endif // !defined(XGBOOST_USE_CUDA) /** @@ -413,9 +413,7 @@ class ArrayInterface { public: ArrayInterface() = default; - explicit ArrayInterface(std::map const &array) { - this->Initialize(array); - } + explicit ArrayInterface(std::map const &array) { this->Initialize(array); } explicit ArrayInterface(Json const &array) { if (IsA(array)) { @@ -533,7 +531,7 @@ class ArrayInterface { // Total number of items size_t n{0}; // Whether the memory is c-contiguous - bool is_contiguous {false}; + bool is_contiguous{false}; // RTTI, initialized to the f16 to avoid masking potential bugs in initialization. ArrayInterfaceHandler::Type type{ArrayInterfaceHandler::kF16}; }; From 167ff7b31cf016150e27ba5f9fc1b0e6ce96f629 Mon Sep 17 00:00:00 2001 From: fis Date: Mon, 15 Nov 2021 18:21:31 +0800 Subject: [PATCH 7/7] Empty dataset. --- src/data/data.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/data/data.cu b/src/data/data.cu index a9e333f31db1..8e13db9ce751 100644 --- a/src/data/data.cu +++ b/src/data/data.cu @@ -28,24 +28,24 @@ template void CopyTensorInfoImpl(Json arr_interface, linalg::Tensor* p_out) { ArrayInterface array(arr_interface); if (array.n == 0) { + p_out->SetDevice(0); return; } CHECK(array.valid.Size() == 0) << "Meta info like label or weight can not have missing value."; auto ptr_device = SetDeviceToPtr(array.data); + p_out->SetDevice(ptr_device); if (array.is_contiguous && array.type == ToDType::kType) { p_out->ModifyInplace([&](HostDeviceVector* data, common::Span shape) { // set shape std::copy(array.shape, array.shape + D, shape.data()); // set data - data->SetDevice(ptr_device); data->Resize(array.n); dh::safe_cuda(cudaMemcpyAsync(data->DevicePointer(), array.data, array.n * sizeof(T), cudaMemcpyDefault)); }); return; } - p_out->SetDevice(ptr_device); p_out->Reshape(array.shape); auto t = p_out->View(ptr_device); linalg::ElementWiseKernelDevice(t, [=] __device__(size_t i, T) {