Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Extend array interface to handle ndarray. (#7434)

Browse Source 
* 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.
This commit is contained in:
Jiaming Yuan 2021-11-16 09:52:15 +08:00 committed by GitHub
parent e27f543deb
commit 55ee272ea8
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
18 changed files with 654 additions and 456 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
include/xgboost/data.h
View File

@ -11,12 +11,14 @@
#include <dmlc/data.h>
#include <dmlc/serializer.h>
#include <xgboost/base.h>
#include <xgboost/span.h>
#include <xgboost/host_device_vector.h>
#include <xgboost/linalg.h>
#include <xgboost/span.h>
#include <xgboost/string_view.h>
#include <algorithm>
#include <memory>
#include <numeric>
#include <algorithm>
#include <string>
#include <utility>
#include <vector>
@ -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;

30
jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
View File

@ -35,13 +35,12 @@ template <typename T> T CheckJvmCall(T const &v, JNIEnv *jenv) {
}
template <typename VCont>
void CopyColumnMask(xgboost::ArrayInterface const &interface,
void CopyColumnMask(xgboost::ArrayInterface<1> const &interface,
std::vector<Json> 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("<t1");
out["mask"]["version"] = Integer(1);
out["mask"]["version"] = Integer(3);
}
template <typename DCont, typename VCont>
void CopyInterface(std::vector<xgboost::ArrayInterface> &interface_arr,
void CopyInterface(std::vector<xgboost::ArrayInterface<1>> &interface_arr,
std::vector<Json> const &columns, cudaMemcpyKind kind,
std::vector<DCont> *p_data, std::vector<VCont> *p_mask,
std::vector<xgboost::Json> *p_out, cudaStream_t stream) {
@ -81,7 +80,7 @@ void CopyInterface(std::vector<xgboost::ArrayInterface> &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<xgboost::ArrayInterface> &interface_arr,
Json{Boolean{false}}};
out["data"] = Array(std::move(j_data));
out["shape"] = Array(std::vector<Json>{Json(Integer(interface.num_rows)),
Json(Integer(interface.num_cols))});
out["shape"] = Array(std::vector<Json>{Json(Integer(interface.Shape(0)))});
if (interface.valid.Data()) {
CopyColumnMask(interface, columns, kind, c, &mask, &out, stream);
}
out["typestr"] = String("<f4");
out["version"] = Integer(1);
out["version"] = Integer(3);
}
}
@ -110,10 +108,10 @@ void CopyMetaInfo(Json *p_interface, dh::device_vector<float> *out, cudaStream_t
auto &j_interface = *p_interface;
CHECK_EQ(get<Array const>(j_interface).size(), 1);
auto object = get<Object>(get<Array>(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<Integer::Int>(RawPtr(*out));
@ -285,11 +283,11 @@ class DataIteratorProxy {
Json features = json_interface["features_str"];
auto json_columns = get<Array const>(features);
std::vector<ArrayInterface> interfaces;
std::vector<ArrayInterface<1>> interfaces;
// Stage the data
for (auto &json_col : json_columns) {
auto column = ArrayInterface(get<Object const>(json_col));
auto column = ArrayInterface<1>(get<Object const>(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<ArrayInterface> in;
std::vector<ArrayInterface<1>> in;
for (auto interface : json_interface) {
auto column = ArrayInterface(get<Object const>(interface));
auto column = ArrayInterface<1>(get<Object const>(interface));
in.emplace_back(column);
}
std::vector<Json> out;

2
src/common/device_helpers.cuh
View File

@ -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));

80
src/data/adapter.h
View File

@ -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<ArrayAdapterBatch> {
public:
explicit ArrayAdapter(StringView array_interface) {
auto j = Json::Load(array_interface);
array_interface_ = ArrayInterface(get<Object const>(j));
array_interface_ = ArrayInterface<2>(get<Object const>(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<size_t>(offset_ + idx, 0),
values_.GetElement(offset_ + idx, 0)};
return {ridx_, TypedIndex<size_t, 1>{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<size_t>(idx, 0);
auto end_offset = indptr_.GetElement<size_t>(idx + 1, 0);
auto begin_no_stride = TypedIndex<size_t, 1>{indptr_}(idx);
auto end_no_stride = TypedIndex<size_t, 1>{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<CSRArrayAdapterBatch>
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<CSRArrayAdapterBatch>
private:
CSRArrayAdapterBatch batch_;
ArrayInterface indptr_;
ArrayInterface indices_;
ArrayInterface values_;
ArrayInterface<1> indptr_;
ArrayInterface<1> indices_;
ArrayInterface<1> values_;
size_t num_cols_;
};

37
src/data/array_interface.cu
View File

@ -1,13 +1,21 @@
/*!
* Copyright 2021 by Contributors
*/
#include "array_interface.h"
#include "../common/common.h"
#include "array_interface.h"
namespace xgboost {
void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) {
switch (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
@ -18,4 +26,31 @@ void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) {
dh::safe_cuda(cudaStreamSynchronize(reinterpret_cast<cudaStream_t>(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

493
src/data/array_interface.h
View File

@ -13,24 +13,23 @@
#include <utility>
#include <vector>
#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 <endian><type><size of type in bytes>.";
}
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 ";
@ -92,49 +91,39 @@ 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:
template <typename T>
static constexpr char TypeChar() {
return
(std::is_floating_point<T>::value ? 'f' :
(std::is_integral<T>::value ?
(std::is_signed<T>::value ? 'i' : 'u') : '\0'));
}
enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 };
template <typename PtrType>
static PtrType GetPtrFromArrayData(std::map<std::string, Json> const& obj) {
if (obj.find("data") == obj.cend()) {
static PtrType GetPtrFromArrayData(std::map<std::string, Json> const &obj) {
auto data_it = obj.find("data");
if (data_it == obj.cend()) {
LOG(FATAL) << "Empty data passed in.";
}
auto p_data = reinterpret_cast<PtrType>(static_cast<size_t>(
get<Integer const>(
get<Array const>(
obj.at("data"))
.at(0))));
auto p_data = reinterpret_cast<PtrType>(
static_cast<size_t>(get<Integer const>(get<Array const>(data_it->second).at(0))));
return p_data;
}
static void Validate(std::map<std::string, Json> const& array) {
static void Validate(std::map<std::string, Json> 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<Null>(stream_it->second)) {
// is cuda, check the version.
if (get<Integer const>(version_it->second) > 3) {
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<String const>(array.at("typestr"));
auto typestr = get<String const>(typestr_it->second);
CHECK(typestr.size() == 3 || typestr.size() == 4) << ArrayInterfaceErrors::TypestrFormat();
if (array.find("shape") == array.cend()) {
@ -149,12 +138,12 @@ class ArrayInterfaceHandler {
// Mask object is also an array interface, but with different requirements.
static size_t ExtractMask(std::map<std::string, Json> const &column,
common::Span<RBitField8::value_type> *p_out) {
auto& s_mask = *p_out;
auto &s_mask = *p_out;
if (column.find("mask") != column.cend()) {
auto const& j_mask = get<Object const>(column.at("mask"));
auto const &j_mask = get<Object const>(column.at("mask"));
Validate(j_mask);
auto p_mask = GetPtrFromArrayData<RBitField8::value_type*>(j_mask);
auto p_mask = GetPtrFromArrayData<RBitField8::value_type *>(j_mask);
auto j_shape = get<Array const>(j_mask.at("shape"));
CHECK_EQ(j_shape.size(), 1) << ArrayInterfaceErrors::Dimension(1);
@ -187,7 +176,7 @@ class ArrayInterfaceHandler {
if (j_mask.find("strides") != j_mask.cend()) {
auto strides = get<Array const>(column.at("strides"));
CHECK_EQ(strides.size(), 1) << ArrayInterfaceErrors::Dimension(1);
CHECK_EQ(get<Integer>(strides.at(0)), type_length) << ArrayInterfaceErrors::Contigious();
CHECK_EQ(get<Integer>(strides.at(0)), type_length) << ArrayInterfaceErrors::Contiguous();
}
s_mask = {p_mask, span_size};
@ -195,77 +184,212 @@ class ArrayInterfaceHandler {
}
return 0;
}
static std::pair<bst_row_t, bst_feature_t> ExtractShape(
std::map<std::string, Json> const& column) {
auto j_shape = get<Array const>(column.at("shape"));
auto typestr = get<String const>(column.at("typestr"));
if (j_shape.size() == 1) {
return {static_cast<bst_row_t>(get<Integer const>(j_shape.at(0))), 1};
} else {
CHECK_EQ(j_shape.size(), 2) << "Only 1-D and 2-D arrays are supported.";
return {static_cast<bst_row_t>(get<Integer const>(j_shape.at(0))),
static_cast<bst_feature_t>(get<Integer const>(j_shape.at(1)))};
/**
* \brief Handle vector inputs. For higher dimension, we require strictly correct shape.
*/
template <int32_t D>
static void HandleRowVector(std::vector<size_t> const &shape, std::vector<size_t> *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<std::string, Json> 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<Null>(strides_it->second)) {
// default strides
*stride_r = cols;
*stride_c = 1;
} else {
// strides specified by the array interface
template <int32_t D>
static void ExtractShape(std::map<std::string, Json> const &array, size_t (&out_shape)[D]) {
auto const &j_shape = get<Array const>(array.at("shape"));
std::vector<size_t> shape_arr(j_shape.size(), 0);
std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(),
[](Json in) { return get<Integer const>(in); });
// handle column vector vs. row vector
HandleRowVector<D>(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);
}
/**
* \brief Extracts the optiona `strides' field and returns whether the array is c-contiguous.
*/
template <int32_t D>
static bool ExtractStride(std::map<std::string, Json> 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<Null>(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 const>(array.at("shape"));
std::vector<size_t> shape_arr(j_shape.size(), 0);
std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(),
[](Json in) { return get<Integer const>(in); });
// Get stride
auto const &j_strides = get<Array const>(strides_it->second);
CHECK_LE(j_strides.size(), 2) << ArrayInterfaceErrors::Dimension(2);
*stride_r = get<Integer const>(j_strides[0]) / itemsize;
size_t n = 1;
if (j_strides.size() == 2) {
n = get<Integer const>(j_strides[1]) / itemsize;
CHECK_EQ(j_strides.size(), j_shape.size()) << "stride and shape don't match.";
std::vector<size_t> stride_arr(j_strides.size(), 0);
std::transform(j_strides.cbegin(), j_strides.cend(), stride_arr.begin(),
[](Json in) { return get<Integer const>(in); });
// Handle column vector vs. row vector
HandleRowVector<D>(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;
}
*stride_c = n;
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);
}
auto valid = rows * (*stride_r) + cols * (*stride_c) >= (rows * cols);
CHECK(valid) << "Invalid strides in array."
<< " strides: (" << (*stride_r) << "," << (*stride_c)
<< "), shape: (" << rows << ", " << cols << ")";
}
static void* ExtractData(std::map<std::string, Json> const &column,
std::pair<size_t, size_t> shape) {
Validate(column);
void* p_data = ArrayInterfaceHandler::GetPtrFromArrayData<void*>(column);
static void *ExtractData(std::map<std::string, Json> const &array, size_t size) {
Validate(array);
void *p_data = ArrayInterfaceHandler::GetPtrFromArrayData<void *>(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 <typename T, typename E = void>
struct ToDType;
// float
template <>
struct ToDType<float> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4;
};
template <>
struct ToDType<double> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF8;
};
template <typename T>
struct ToDType<T,
std::enable_if_t<std::is_same<T, long double>::value && sizeof(long double) == 16>> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF16;
};
// uint
template <>
struct ToDType<uint8_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU1;
};
template <>
struct ToDType<uint16_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU2;
};
template <>
struct ToDType<uint32_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU4;
};
template <>
struct ToDType<uint64_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU8;
};
// int
template <>
struct ToDType<int8_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI1;
};
template <>
struct ToDType<int16_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI2;
};
template <>
struct ToDType<int32_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI4;
};
template <>
struct ToDType<int64_t> {
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
*
* <a href="https://numpy.org/doc/stable/reference/arrays.interface.html">numpy</a>.
*
* \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 <int32_t D, bool allow_mask = (D == 1)>
class ArrayInterface {
void Initialize(std::map<std::string, Json> 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<std::string, Json> const &array) {
ArrayInterfaceHandler::Validate(array);
auto typestr = get<String const>(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<RBitField8::value_type> s_mask;
size_t n_bits = ArrayInterfaceHandler::ExtractMask(array, &s_mask);
@ -273,18 +397,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. "
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<Null>(stream_it->second)) {
int64_t stream = get<Integer const>(stream_it->second);
@ -292,151 +411,147 @@ 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<std::string, Json> const &array) { this->Initialize(array); }
explicit ArrayInterface(std::map<std::string, Json> const &column,
bool allow_mask = true) {
this->Initialize(column, allow_mask);
}
explicit ArrayInterface(StringView str, bool allow_mask = true) {
auto jinterface = Json::Load(str);
if (IsA<Object>(jinterface)) {
this->Initialize(get<Object const>(jinterface), allow_mask);
explicit ArrayInterface(Json const &array) {
if (IsA<Object>(array)) {
this->Initialize(get<Object const>(array));
return;
}
if (IsA<Array>(jinterface)) {
CHECK_EQ(get<Array const>(jinterface).size(), 1)
if (IsA<Array>(array)) {
CHECK_EQ(get<Array const>(array).size(), 1)
<< "Column: " << ArrayInterfaceErrors::Dimension(1);
this->Initialize(get<Object const>(get<Array const>(jinterface)[0]), allow_mask);
this->Initialize(get<Object const>(get<Array const>(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<size_t>(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<D>{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 <typename Fn>
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<float *>(data));
case kF8:
return func(reinterpret_cast<double *>(data));
case T::kF4:
return func(reinterpret_cast<float const *>(data));
case T::kF8:
return func(reinterpret_cast<double const *>(data));
#ifdef __CUDA_ARCH__
case kF16: {
case T::kF16: {
// CUDA device code doesn't support long double.
SPAN_CHECK(false);
return func(reinterpret_cast<double *>(data));
return func(reinterpret_cast<double const *>(data));
}
#else
case kF16:
return func(reinterpret_cast<long double *>(data));
case T::kF16:
return func(reinterpret_cast<long double const *>(data));
#endif
case kI1:
return func(reinterpret_cast<int8_t *>(data));
case kI2:
return func(reinterpret_cast<int16_t *>(data));
case kI4:
return func(reinterpret_cast<int32_t *>(data));
case kI8:
return func(reinterpret_cast<int64_t *>(data));
case kU1:
return func(reinterpret_cast<uint8_t *>(data));
case kU2:
return func(reinterpret_cast<uint16_t *>(data));
case kU4:
return func(reinterpret_cast<uint32_t *>(data));
case kU8:
return func(reinterpret_cast<uint64_t *>(data));
case T::kI1:
return func(reinterpret_cast<int8_t const *>(data));
case T::kI2:
return func(reinterpret_cast<int16_t const *>(data));
case T::kI4:
return func(reinterpret_cast<int32_t const *>(data));
case T::kI8:
return func(reinterpret_cast<int64_t const *>(data));
case T::kU1:
return func(reinterpret_cast<uint8_t const *>(data));
case T::kU2:
return func(reinterpret_cast<uint16_t const *>(data));
case T::kU4:
return func(reinterpret_cast<uint32_t const *>(data));
case T::kU8:
return func(reinterpret_cast<uint64_t const *>(data));
}
SPAN_CHECK(false);
return func(reinterpret_cast<uint64_t *>(data));
return func(reinterpret_cast<uint64_t const *>(data));
}
XGBOOST_DEVICE size_t ElementSize() {
return this->DispatchCall([](auto* p_values) {
return sizeof(std::remove_pointer_t<decltype(p_values)>);
XGBOOST_DEVICE size_t constexpr ElementSize() {
return this->DispatchCall(
[](auto *p_values) { return sizeof(std::remove_pointer_t<decltype(p_values)>); });
}
template <typename T = float, typename... Index>
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<T>(p_values[offset]);
});
}
template <typename T = float>
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]; });
}
// 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 const *data{nullptr};
// Total number of items
size_t n{0};
// Whether the memory is c-contiguous
bool is_contiguous{false};
// RTTI, initialized to the f16 to avoid masking potential bugs in initialization.
ArrayInterfaceHandler::Type type{ArrayInterfaceHandler::kF16};
};
template <typename T> std::string MakeArrayInterface(T const *data, size_t n) {
Json arr{Object{}};
arr["data"] = Array(std::vector<Json>{
Json{Integer{reinterpret_cast<int64_t>(data)}}, Json{Boolean{false}}});
arr["shape"] = Array{std::vector<Json>{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 <typename T, int32_t D>
struct TypedIndex {
ArrayInterface<D> const &array;
template <typename... I>
XGBOOST_DEVICE T operator()(I &&...ind) const {
static_assert(sizeof...(ind) <= D, "Invalid index.");
return array.template operator()<T>(ind...);
}
typestr.push_back(ArrayInterfaceHandler::TypeChar<T>());
typestr += std::to_string(sizeof(T));
arr["typestr"] = typestr;
arr["version"] = 3;
std::string str;
Json::Dump(arr, &str);
return str;
};
template <int32_t D>
inline void CheckArrayInterface(StringView key, ArrayInterface<D> const &array) {
CHECK(!array.valid.Data()) << "Meta info " << key << " should be dense, found validity mask";
}
} // namespace xgboost
#endif // XGBOOST_DATA_ARRAY_INTERFACE_H_

18
src/data/data.cc
View File

@ -1,5 +1,5 @@
/*!
* Copyright 2015-2020 by Contributors
* Copyright 2015-2021 by Contributors
* \file data.cc
*/
#include <dmlc/registry.h>
@ -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<bst_group_t> 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<uint32_t> 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)

184
src/data/data.cu
View File

@ -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<float>* out) {
auto SetDeviceToPtr = [](void* ptr) {
namespace {
auto SetDeviceToPtr(void const* 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) {
template <typename T, int32_t D>
void CopyTensorInfoImpl(Json arr_interface, linalg::Tensor<T, D>* p_out) {
ArrayInterface<D> array(arr_interface);
if (array.n == 0) {
p_out->SetDevice(0);
return;
}
out->SetDevice(ptr_device);
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);
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);
if (array.is_contiguous && array.type == ToDType<T>::kType) {
p_out->ModifyInplace([&](HostDeviceVector<T>* data, common::Span<size_t, D> shape) {
// set shape
std::copy(array.shape, array.shape + D, shape.data());
// set data
data->Resize(array.n);
dh::safe_cuda(cudaMemcpyAsync(data->DevicePointer(), array.data, array.n * sizeof(T),
cudaMemcpyDefault));
});
return;
}
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<T, D>{array}, linalg::UnravelIndex<D>(i, array.shape));
});
}
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;
}
} // anonymous namespace
void CopyGroupInfoImpl(ArrayInterface column, std::vector<bst_group_t>* out) {
CHECK(column.type != ArrayInterface::kF4 && column.type != ArrayInterface::kF8)
void CopyGroupInfoImpl(ArrayInterface<1> column, std::vector<bst_group_t>* 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<bst_group_t> temp(column.num_rows);
auto d_tmp = temp.data();
dh::TemporaryArray<bst_group_t> temp(column.Shape(0));
auto d_tmp = temp.data().get();
dh::LaunchN(column.num_rows, [=] __device__(size_t idx) {
d_tmp[idx] = column.GetElement<size_t>(idx, 0);
});
auto length = column.num_rows;
dh::LaunchN(column.Shape(0),
[=] __device__(size_t idx) { d_tmp[idx] = TypedIndex<size_t, 1>{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<bst_group_t> *p_group_ptr) {
void CopyQidImpl(ArrayInterface<1> array_interface, std::vector<bst_group_t>* p_group_ptr) {
auto &group_ptr_ = *p_group_ptr;
auto it = dh::MakeTransformIterator<uint32_t>(
thrust::make_counting_iterator(0ul),
[array_interface] __device__(size_t i) {
return array_interface.GetElement<uint32_t>(i, 0);
thrust::make_counting_iterator(0ul), [array_interface] __device__(size_t i) {
return TypedIndex<uint32_t, 1>{array_interface}(i);
});
dh::caching_device_vector<bool> 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<uint32_t>(i, 0) >
array_interface.GetElement<uint32_t>(i + 1, 0)) {
dh::LaunchN(array_interface.Shape(0) - 1, [=] __device__(size_t i) {
auto typed = TypedIndex<uint32_t, 1>{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<uint32_t> out(array_interface.num_rows);
dh::caching_device_vector<uint32_t> cnt(array_interface.num_rows);
dh::caching_device_vector<uint32_t> out(array_interface.Shape(0));
dh::caching_device_vector<uint32_t> cnt(array_interface.Shape(0));
HostDeviceVector<int> 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<char> 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<bst_group_t> 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<float, 3> 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();
}
if (key == "label") {
CopyInfoImpl(array_interface, &labels_);
auto ptr = labels_.ConstDevicePointer();
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") {
CopyInfoImpl(array_interface, &weights_);
auto ptr = weights_.ConstDevicePointer();
auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(),
WeightsCheck{});
CHECK(valid) << "Weights must be positive values.";
} else if (key == "group") {
// uint info
if (key == "group") {
auto array_interface{ArrayInterface<1>(array)};
CopyGroupInfoImpl(array_interface, &group_ptr_);
ValidateQueryGroup(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<float, 1> t;
CopyTensorInfoImpl(array, &t);
if (key == "label") {
this->labels_ = std::move(*t.Data());
auto ptr = labels_.ConstDevicePointer();
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") {
this->weights_ = std::move(*t.Data());
auto ptr = weights_.ConstDevicePointer();
auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(), data::WeightsCheck{});
CHECK(valid) << "Weights must be positive values.";
} 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;
}
}

47
src/data/device_adapter.cuh
View File

@ -20,7 +20,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo {
public:
CudfAdapterBatch() = default;
CudfAdapterBatch(common::Span<ArrayInterface> columns, size_t num_rows)
CudfAdapterBatch(common::Span<ArrayInterface<1>> 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<float>::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<ArrayInterface> columns_;
common::Span<ArrayInterface<1>> columns_;
size_t num_rows_;
};
@ -101,9 +101,9 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
auto const& typestr = get<String const>(json_columns[0]["typestr"]);
CHECK_EQ(typestr.size(), 3) << ArrayInterfaceErrors::TypestrFormat();
std::vector<ArrayInterface> columns;
auto first_column = ArrayInterface(get<Object const>(json_columns[0]));
num_rows_ = first_column.num_rows;
std::vector<ArrayInterface<1>> columns;
auto first_column = ArrayInterface<1>(get<Object const>(json_columns[0]));
num_rows_ = first_column.Shape(0);
if (num_rows_ == 0) {
return;
}
@ -112,13 +112,12 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
CHECK_NE(device_idx_, -1);
dh::safe_cuda(cudaSetDevice(device_idx_));
for (auto& json_col : json_columns) {
auto column = ArrayInterface(get<Object const>(json_col));
auto column = ArrayInterface<1>(get<Object const>(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<CudfAdapterBatch> {
private:
CudfAdapterBatch batch_;
dh::device_vector<ArrayInterface> columns_;
dh::device_vector<ArrayInterface<1>> columns_;
size_t num_rows_{0};
int device_idx_;
};
@ -143,23 +142,23 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
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<CupyAdapterBatch> {
@ -167,9 +166,9 @@ class CupyAdapter : public detail::SingleBatchDataIter<CupyAdapterBatch> {
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<Object const>(json_array_interface), false);
array_interface_ = ArrayInterface<2>(get<Object const>(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<CupyAdapterBatch> {
}
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};
};

13
src/data/file_iterator.h
View File

@ -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,15 +59,13 @@ 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,
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

40
src/data/validation.h Normal file
View File

@ -0,0 +1,40 @@
/*!
* Copyright 2021 by XGBoost Contributors
*/
#ifndef XGBOOST_DATA_VALIDATION_H_
#define XGBOOST_DATA_VALIDATION_H_
#include <cmath>
#include <vector>
#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<bst_group_t> 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_

7
tests/cpp/data/test_adapter.cc
View File

@ -41,9 +41,10 @@ TEST(Adapter, CSRArrayAdapter) {
HostDeviceVector<bst_feature_t> 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()},

66
tests/cpp/data/test_array_interface.cc
View File

@ -11,21 +11,22 @@ TEST(ArrayInterface, Initialize) {
size_t constexpr kRows = 10, kCols = 10;
HostDeviceVector<float> 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<size_t> 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<size_t const, 2>{
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<Object>(column);
std::pair<size_t, size_t> shape{kRows, kCols};
std::string typestr{"<f4"};
size_t n = kRows * kCols;
// missing version
EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
dmlc::Error);
column["version"] = Integer(static_cast<Integer::Int>(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("<f4");
// nullptr is not valid
EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n),
dmlc::Error);
HostDeviceVector<float> storage;
@ -63,22 +63,52 @@ TEST(ArrayInterface, Error) {
Json(Integer(reinterpret_cast<Integer::Int>(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<float> 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<float> 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<float>::kType == ArrayInterfaceHandler::kF4, "");
static_assert(ToDType<double>::kType == ArrayInterfaceHandler::kF8, "");
static_assert(ToDType<uint32_t>::kType == ArrayInterfaceHandler::kU4, "");
static_assert(ToDType<uint64_t>::kType == ArrayInterfaceHandler::kU8, "");
static_assert(ToDType<int32_t>::kType == ArrayInterfaceHandler::kI4, "");
static_assert(ToDType<int64_t>::kType == ArrayInterfaceHandler::kI8, "");
}
} // namespace xgboost

15
tests/cpp/data/test_array_interface.cu
View File

@ -32,11 +32,24 @@ TEST(ArrayInterface, Stream) {
dh::caching_device_vector<uint64_t> 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<float> h_data(10);
ASSERT_FALSE(ArrayInterfaceHandler::IsCudaPtr(h_data.data()));
dh::safe_cuda(cudaGetLastError());
dh::device_vector<float> 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

10
tests/cpp/data/test_array_interface.h
View File

@ -19,6 +19,7 @@ Json GenerateDenseColumn(std::string const& typestr, size_t kRows,
std::vector<Json> j_shape {Json(Integer(static_cast<Integer::Int>(kRows)))};
column["shape"] = Array(j_shape);
column["strides"] = Array(std::vector<Json>{Json(Integer(static_cast<Integer::Int>(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<Integer::Int>(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<Json> j_shape {Json(Integer(static_cast<Integer::Int>(kRows)))};
column["shape"] = Array(j_shape);
column["strides"] = Array(std::vector<Json>{Json(Integer(static_cast<Integer::Int>(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<Integer::Int>(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<Integer::Int>(data.data().get()))),
Json(Boolean(false))};
array_interface["data"] = j_data;
array_interface["version"] = Integer(static_cast<Integer::Int>(1));
array_interface["version"] = 3;
array_interface["typestr"] = String(typestr);
array_interface["stream"] = nullptr;
return array_interface;
}
} // namespace xgboost

8
tests/cpp/data/test_iterative_device_dmatrix.cu
View File

@ -103,9 +103,9 @@ TEST(IterativeDeviceDMatrix, RowMajor) {
auto j_interface =
Json::Load({interface_str.c_str(), interface_str.size()});
ArrayInterface loaded {get<Object const>(j_interface)};
ArrayInterface<2> loaded {get<Object const>(j_interface)};
std::vector<float> h_data(cols * rows);
common::Span<float> s_data{static_cast<float*>(loaded.data), cols * rows};
common::Span<float const> s_data{static_cast<float const*>(loaded.data), cols * rows};
dh::CopyDeviceSpanToVector(&h_data, s_data);
for(auto i = 0ull; i < rows * cols; i++) {
@ -128,9 +128,9 @@ TEST(IterativeDeviceDMatrix, RowMajorMissing) {
std::string interface_str = iter.AsArray();
auto j_interface =
Json::Load({interface_str.c_str(), interface_str.size()});
ArrayInterface loaded {get<Object const>(j_interface)};
ArrayInterface<2> loaded {get<Object const>(j_interface)};
std::vector<float> h_data(cols * rows);
common::Span<float> s_data{static_cast<float*>(loaded.data), cols * rows};
common::Span<float const> s_data{static_cast<float const*>(loaded.data), cols * rows};
dh::CopyDeviceSpanToVector(&h_data, s_data);
h_data[1] = kMissing;
h_data[5] = kMissing;

2
tests/cpp/data/test_metainfo.cu
View File

@ -1,4 +1,4 @@
/*! Copyright 2019 by Contributors */
/*! Copyright 2019-2021 by XGBoost Contributors */
#include <gtest/gtest.h>
#include <xgboost/data.h>

2
tests/cpp/helpers.h
View File

@ -198,7 +198,7 @@ Json GetArrayInterface(HostDeviceVector<T> *storage, size_t rows, size_t cols) {
array_interface["shape"][0] = rows;
array_interface["shape"][1] = cols;
char t = ArrayInterfaceHandler::TypeChar<T>();
char t = linalg::detail::ArrayInterfaceHandler::TypeChar<T>();
array_interface["typestr"] = String(std::string{"<"} + t + std::to_string(sizeof(T)));
array_interface["version"] = 1;
return array_interface;