xgboost

linalg.h
Go to the documentation of this file.
1
6 #ifndef XGBOOST_LINALG_H_
7 #define XGBOOST_LINALG_H_
8
9 #include <dmlc/endian.h>
10 #include <xgboost/base.h>
11 #include <xgboost/host_device_vector.h>
12 #include <xgboost/json.h>
13 #include <xgboost/span.h>
14
15 #include <algorithm>
16 #include <cassert>
17 #include <limits>
18 #include <string>
19 #include <type_traits>
20 #include <utility>
21 #include <vector>
22
23 // decouple it from xgboost.
24 #ifndef LINALG_HD
25 #if defined(__CUDA__) || defined(__NVCC__)
26 #define LINALG_HD __host__ __device__
27 #else
28 #define LINALG_HD
29 #endif // defined (__CUDA__) || defined(__NVCC__)
30 #endif // LINALG_HD
31
32 namespace xgboost {
33 namespace linalg {
34 namespace detail {
35
36 struct ArrayInterfaceHandler {
37 template <typename T>
38 static constexpr char TypeChar() {
39 return (std::is_floating_point<T>::value
40 ? 'f'
41 : (std::is_integral<T>::value ? (std::is_signed<T>::value ? 'i' : 'u') : '\0'));
42 }
43 };
44
45 template <size_t dim, typename S, typename Head, size_t D>
46 constexpr size_t Offset(S (&strides)[D], size_t n, Head head) {
47 static_assert(dim < D, "");
48 return n + head * strides[dim];
49 }
50
51 template <size_t dim, typename S, size_t D, typename Head, typename... Tail>
52 constexpr std::enable_if_t<sizeof...(Tail) != 0, size_t> Offset(S (&strides)[D], size_t n,
53 Head head, Tail &&...rest) {
54 static_assert(dim < D, "");
55 return Offset<dim + 1>(strides, n + (head * strides[dim]), std::forward<Tail>(rest)...);
56 }
57
58 template <int32_t D, bool f_array = false>
59 constexpr void CalcStride(size_t const (&shape)[D], size_t (&stride)[D]) {
60 if (f_array) {
61 stride[0] = 1;
62 for (int32_t s = 1; s < D; ++s) {
63 stride[s] = shape[s - 1] * stride[s - 1];
64 }
65 } else {
66 stride[D - 1] = 1;
67 for (int32_t s = D - 2; s >= 0; --s) {
68 stride[s] = shape[s + 1] * stride[s + 1];
69 }
70 }
71 }
72
73 struct AllTag {};
74
75 struct IntTag {};
76
77 template <typename I>
78 struct RangeTag {
79 I beg;
80 I end;
81 constexpr size_t Size() const { return end - beg; }
82 };
83
87 template <typename T>
88 constexpr int32_t CalcSliceDim() {
89 return std::is_same<T, IntTag>::value ? 0 : 1;
90 }
91
92 template <typename T, typename... S>
93 constexpr std::enable_if_t<sizeof...(S) != 0, int32_t> CalcSliceDim() {
94 return CalcSliceDim<T>() + CalcSliceDim<S...>();
95 }
96
97 template <int32_t D>
98 constexpr size_t CalcSize(size_t (&shape)[D]) {
99 size_t size = 1;
100 for (auto d : shape) {
101 size *= d;
102 }
103 return size;
104 }
105
106 template <typename S>
107 using RemoveCRType = std::remove_const_t<std::remove_reference_t<S>>;
108
109 template <typename S>
110 using IndexToTag = std::conditional_t<std::is_integral<RemoveCRType<S>>::value, IntTag, S>;
111
112 template <int32_t n, typename Fn>
113 LINALG_HD constexpr auto UnrollLoop(Fn fn) {
114 #if defined __CUDA_ARCH__
115 #pragma unroll n
116 #endif // defined __CUDA_ARCH__
117 for (int32_t i = 0; i < n; ++i) {
118 fn(i);
119 }
120 }
121
122 template <typename T>
123 int32_t NativePopc(T v) {
124 int c = 0;
125 for (; v != 0; v &= v - 1) c++;
126 return c;
127 }
128
129 inline LINALG_HD int Popc(uint32_t v) {
130 #if defined(__CUDA_ARCH__)
131 return __popc(v);
132 #elif defined(__GNUC__) || defined(__clang__)
133 return __builtin_popcount(v);
134 #elif defined(_MSC_VER)
135 return __popcnt(v);
136 #else
137 return NativePopc(v);
138 #endif // compiler
139 }
140
141 inline LINALG_HD int Popc(uint64_t v) {
142 #if defined(__CUDA_ARCH__)
143 return __popcll(v);
144 #elif defined(__GNUC__) || defined(__clang__)
145 return __builtin_popcountll(v);
146 #elif defined(_MSC_VER)
147 return __popcnt64(v);
148 #else
149 return NativePopc(v);
150 #endif // compiler
151 }
152
153 template <class T, std::size_t N, std::size_t... Idx>
154 constexpr auto Arr2Tup(T (&arr)[N], std::index_sequence<Idx...>) {
155 return std::make_tuple(arr[Idx]...);
156 }
157
158 template <class T, std::size_t N>
159 constexpr auto Arr2Tup(T (&arr)[N]) {
160 return Arr2Tup(arr, std::make_index_sequence<N>{});
161 }
162
163 // uint division optimization inspired by the CIndexer in cupy. Division operation is
164 // slow on both CPU and GPU, especially 64 bit integer. So here we first try to avoid 64
165 // bit when the index is smaller, then try to avoid division when it's exp of 2.
166 template <typename I, int32_t D>
167 LINALG_HD auto UnravelImpl(I idx, common::Span<size_t const, D> shape) {
168 size_t index[D]{0};
169 static_assert(std::is_signed<decltype(D)>::value,
170 "Don't change the type without changing the for loop.");
171 for (int32_t dim = D; --dim > 0;) {
172 auto s = static_cast<std::remove_const_t<std::remove_reference_t<I>>>(shape[dim]);
173 if (s & (s - 1)) {
174 auto t = idx / s;
175 index[dim] = idx - t * s;
176 idx = t;
177 } else { // exp of 2
178 index[dim] = idx & (s - 1);
179 idx >>= Popc(s - 1);
180 }
181 }
182 index[0] = idx;
183 return Arr2Tup(index);
184 }
185
186 template <size_t dim, typename I, int32_t D>
187 void ReshapeImpl(size_t (&out_shape)[D], I s) {
188 static_assert(dim < D, "");
189 out_shape[dim] = s;
190 }
191
192 template <size_t dim, int32_t D, typename... S, typename I,
193 std::enable_if_t<sizeof...(S) != 0> * = nullptr>
194 void ReshapeImpl(size_t (&out_shape)[D], I &&s, S &&...rest) {
195 static_assert(dim < D, "");
196 out_shape[dim] = s;
197 ReshapeImpl<dim + 1>(out_shape, std::forward<S>(rest)...);
198 }
199
200 template <typename Fn, typename Tup, size_t... I>
201 LINALG_HD decltype(auto) constexpr Apply(Fn &&f, Tup &&t, std::index_sequence<I...>) {
202 return f(std::get<I>(t)...);
203 }
204
211 template <typename Fn, typename Tup>
212 LINALG_HD decltype(auto) constexpr Apply(Fn &&f, Tup &&t) {
213 constexpr auto kSize = std::tuple_size<Tup>::value;
214 return Apply(std::forward<Fn>(f), std::forward<Tup>(t), std::make_index_sequence<kSize>{});
215 }
216 } // namespace detail
217
221 constexpr detail::AllTag All() { return {}; }
225 template <typename I>
226 constexpr detail::RangeTag<I> Range(I beg, I end) {
227 return {beg, end};
228 }
229
243 template <typename T, int32_t kDim>
244 class TensorView {
245 public:
246 using ShapeT = size_t[kDim];
247 using StrideT = ShapeT;
248
249 private:
250 StrideT stride_{1};
251 ShapeT shape_{0};
252 common::Span<T> data_;
253 T *ptr_{nullptr}; // pointer of data_ to avoid bound check.
254
255 size_t size_{0};
256 int32_t device_{-1};
257
258 // Unlike `Tensor`, the data_ can have arbitrary size since this is just a view.
259 LINALG_HD void CalcSize() {
260 if (data_.empty()) {
261 size_ = 0;
262 } else {
263 size_ = detail::CalcSize(shape_);
264 }
265 }
266
267 template <size_t old_dim, size_t new_dim, int32_t D, typename I>
268 LINALG_HD size_t MakeSliceDim(size_t new_shape[D], size_t new_stride[D],
269 detail::RangeTag<I> &&range) const {
270 static_assert(new_dim < D, "");
271 static_assert(old_dim < kDim, "");
272 new_stride[new_dim] = stride_[old_dim];
273 new_shape[new_dim] = range.Size();
274 assert(static_cast<decltype(shape_[old_dim])>(range.end) <= shape_[old_dim]);
275
276 auto offset = stride_[old_dim] * range.beg;
277 return offset;
278 }
282 template <size_t old_dim, size_t new_dim, int32_t D, typename I, typename... S>
283 LINALG_HD size_t MakeSliceDim(size_t new_shape[D], size_t new_stride[D],
284 detail::RangeTag<I> &&range, S &&...slices) const {
285 static_assert(new_dim < D, "");
286 static_assert(old_dim < kDim, "");
287 new_stride[new_dim] = stride_[old_dim];
288 new_shape[new_dim] = range.Size();
289 assert(static_cast<decltype(shape_[old_dim])>(range.end) <= shape_[old_dim]);
290
291 auto offset = stride_[old_dim] * range.beg;
292 return MakeSliceDim<old_dim + 1, new_dim + 1, D>(new_shape, new_stride,
293 std::forward<S>(slices)...) +
294 offset;
295 }
296
297 template <size_t old_dim, size_t new_dim, int32_t D>
298 LINALG_HD size_t MakeSliceDim(size_t new_shape[D], size_t new_stride[D], detail::AllTag) const {
299 static_assert(new_dim < D, "");
300 static_assert(old_dim < kDim, "");
301 new_stride[new_dim] = stride_[old_dim];
302 new_shape[new_dim] = shape_[old_dim];
303 return 0;
304 }
308 template <size_t old_dim, size_t new_dim, int32_t D, typename... S>
309 LINALG_HD size_t MakeSliceDim(size_t new_shape[D], size_t new_stride[D], detail::AllTag,
310 S &&...slices) const {
311 static_assert(new_dim < D, "");
312 static_assert(old_dim < kDim, "");
313 new_stride[new_dim] = stride_[old_dim];
314 new_shape[new_dim] = shape_[old_dim];
315 return MakeSliceDim<old_dim + 1, new_dim + 1, D>(new_shape, new_stride,
316 std::forward<S>(slices)...);
317 }
318
319 template <size_t old_dim, size_t new_dim, int32_t D, typename Index>
320 LINALG_HD size_t MakeSliceDim(size_t new_shape[D], size_t new_stride[D], Index i) const {
321 static_assert(old_dim < kDim, "");
322 return stride_[old_dim] * i;
323 }
327 template <size_t old_dim, size_t new_dim, int32_t D, typename Index, typename... S>
328 LINALG_HD std::enable_if_t<std::is_integral<Index>::value, size_t> MakeSliceDim(
329 size_t new_shape[D], size_t new_stride[D], Index i, S &&...slices) const {
330 static_assert(old_dim < kDim, "");
331 auto offset = stride_[old_dim] * i;
332 auto res =
333 MakeSliceDim<old_dim + 1, new_dim, D>(new_shape, new_stride, std::forward<S>(slices)...);
334 return res + offset;
335 }
336
337 public:
338 size_t constexpr static kValueSize = sizeof(T);
339 size_t constexpr static kDimension = kDim;
340
341 public:
353 template <typename I, int32_t D>
354 LINALG_HD TensorView(common::Span<T> data, I const (&shape)[D], int32_t device)
355 : data_{data}, ptr_{data_.data()}, device_{device} {
356 static_assert(D > 0 && D <= kDim, "Invalid shape.");
357 // shape
358 detail::UnrollLoop<D>([&](auto i) { shape_[i] = shape[i]; });
359 for (auto i = D; i < kDim; ++i) {
360 shape_[i] = 1;
361 }
362 // stride
363 detail::CalcStride(shape_, stride_);
364 // size
365 this->CalcSize();
366 }
367
372 template <typename I, int32_t D>
373 LINALG_HD TensorView(common::Span<T> data, I const (&shape)[D], I const (&stride)[D],
374 int32_t device)
375 : data_{data}, ptr_{data_.data()}, device_{device} {
376 static_assert(D == kDim, "Invalid shape & stride.");
377 detail::UnrollLoop<D>([&](auto i) {
378 shape_[i] = shape[i];
379 stride_[i] = stride[i];
380 });
381 this->CalcSize();
382 }
383
384 template <
385 typename U,
386 std::enable_if_t<common::detail::IsAllowedElementTypeConversion<U, T>::value> * = nullptr>
387 LINALG_HD TensorView(TensorView<U, kDim> const &that) // NOLINT
388 : data_{that.Values()}, ptr_{data_.data()}, size_{that.Size()}, device_{that.DeviceIdx()} {
389 detail::UnrollLoop<kDim>([&](auto i) {
390 stride_[i] = that.Stride(i);
391 shape_[i] = that.Shape(i);
392 });
393 }
394
408 template <typename... Index>
409 LINALG_HD T &operator()(Index &&...index) {
410 static_assert(sizeof...(index) <= kDim, "Invalid index.");
411 size_t offset = detail::Offset<0ul>(stride_, 0ul, std::forward<Index>(index)...);
412 assert(offset < data_.size() && "Out of bound access.");
413 return ptr_[offset];
414 }
418 template <typename... Index>
419 LINALG_HD T const &operator()(Index &&...index) const {
420 static_assert(sizeof...(index) <= kDim, "Invalid index.");
421 size_t offset = detail::Offset<0ul>(stride_, 0ul, std::forward<Index>(index)...);
422 assert(offset < data_.size() && "Out of bound access.");
423 return ptr_[offset];
424 }
425
439 template <typename... S>
440 LINALG_HD auto Slice(S &&...slices) const {
441 static_assert(sizeof...(slices) <= kDim, "Invalid slice.");
442 int32_t constexpr kNewDim{detail::CalcSliceDim<detail::IndexToTag<S>...>()};
443 size_t new_shape[kNewDim];
444 size_t new_stride[kNewDim];
445 auto offset = MakeSliceDim<0, 0, kNewDim>(new_shape, new_stride, std::forward<S>(slices)...);
446 // ret is a different type due to changed dimension, so we can not access its private
447 // fields.
448 TensorView<T, kNewDim> ret{data_.subspan(data_.empty() ? 0 : offset), new_shape, new_stride,
449 device_};
450 return ret;
451 }
452
453 LINALG_HD auto Shape() const { return common::Span<size_t const, kDim>{shape_}; }
457 LINALG_HD auto Shape(size_t i) const { return shape_[i]; }
458 LINALG_HD auto Stride() const { return common::Span<size_t const, kDim>{stride_}; }
462 LINALG_HD auto Stride(size_t i) const { return stride_[i]; }
463
467 LINALG_HD size_t Size() const { return size_; }
471 LINALG_HD bool Contiguous() const {
472 return data_.size() == this->Size() || this->CContiguous() || this->FContiguous();
473 }
477 LINALG_HD bool CContiguous() const {
478 StrideT stride;
479 static_assert(std::is_same<decltype(stride), decltype(stride_)>::value, "");
480 // It's contiguous if the stride can be calculated from shape.
481 detail::CalcStride(shape_, stride);
482 return common::Span<size_t const, kDim>{stride_} == common::Span<size_t const, kDim>{stride};
483 }
487 LINALG_HD bool FContiguous() const {
488 StrideT stride;
489 static_assert(std::is_same<decltype(stride), decltype(stride_)>::value, "");
490 // It's contiguous if the stride can be calculated from shape.
491 detail::CalcStride<kDim, true>(shape_, stride);
492 return common::Span<size_t const, kDim>{stride_} == common::Span<size_t const, kDim>{stride};
493 }
497 LINALG_HD auto Values() const -> decltype(data_) const & { return data_; }
501 LINALG_HD auto DeviceIdx() const { return device_; }
502 };
503
507 template <typename Container, typename I, int32_t D,
508 std::enable_if_t<!common::detail::IsSpan<Container>::value> * = nullptr>
509 auto MakeTensorView(Container &data, I const (&shape)[D], int32_t device) { // NOLINT
510 using T = typename Container::value_type;
511 return TensorView<T, D>{data, shape, device};
512 }
513
514 template <typename T, typename I, int32_t D>
515 LINALG_HD auto MakeTensorView(common::Span<T> data, I const (&shape)[D], int32_t device) {
516 return TensorView<T, D>{data, shape, device};
517 }
518
522 template <size_t D>
523 LINALG_HD auto UnravelIndex(size_t idx, common::Span<size_t const, D> shape) {
524 if (idx > std::numeric_limits<uint32_t>::max()) {
525 return detail::UnravelImpl<uint64_t, D>(static_cast<uint64_t>(idx), shape);
526 } else {
527 return detail::UnravelImpl<uint32_t, D>(static_cast<uint32_t>(idx), shape);
528 }
529 }
530
536 template <typename T>
537 using VectorView = TensorView<T, 1>;
538
546 template <typename T>
547 auto MakeVec(T *ptr, size_t s, int32_t device = -1) {
548 return linalg::TensorView<T, 1>{{ptr, s}, {s}, device};
549 }
550
551 template <typename T>
552 auto MakeVec(HostDeviceVector<T> *data) {
553 return MakeVec(data->DeviceIdx() == -1 ? data->HostPointer() : data->DevicePointer(),
554 data->Size(), data->DeviceIdx());
555 }
556
557 template <typename T>
558 auto MakeVec(HostDeviceVector<T> const *data) {
559 return MakeVec(data->DeviceIdx() == -1 ? data->ConstHostPointer() : data->ConstDevicePointer(),
560 data->Size(), data->DeviceIdx());
561 }
562
568 template <typename T>
569 using MatrixView = TensorView<T, 2>;
570
577 template <typename T, int32_t D>
578 Json ArrayInterface(TensorView<T const, D> const &t) {
579 Json array_interface{Object{}};
580 array_interface["data"] = std::vector<Json>(2);
581 array_interface["data"][0] = Integer{reinterpret_cast<int64_t>(t.Values().data())};
582 array_interface["data"][1] = Boolean{true};
583 if (t.DeviceIdx() >= 0) {
584 // Change this once we have different CUDA stream.
585 array_interface["stream"] = Null{};
586 }
587 std::vector<Json> shape(t.Shape().size());
588 std::vector<Json> stride(t.Stride().size());
589 for (size_t i = 0; i < t.Shape().size(); ++i) {
590 shape[i] = Integer(t.Shape(i));
591 stride[i] = Integer(t.Stride(i) * sizeof(T));
592 }
593 array_interface["shape"] = Array{shape};
594 array_interface["strides"] = Array{stride};
595 array_interface["version"] = 3;
596
597 char constexpr kT = detail::ArrayInterfaceHandler::TypeChar<T>();
598 static_assert(kT != '\0', "");
599 if (DMLC_LITTLE_ENDIAN) {
600 array_interface["typestr"] = String{"<" + (kT + std::to_string(sizeof(T)))};
601 } else {
602 array_interface["typestr"] = String{">" + (kT + std::to_string(sizeof(T)))};
603 }
604 return array_interface;
605 }
606
610 template <typename T, int32_t D>
611 Json ArrayInterface(TensorView<T, D> const &t) {
612 TensorView<T const, D> const &as_const = t;
613 auto res = ArrayInterface(as_const);
614 res["data"][1] = Boolean{false};
615 return res;
616 }
617
621 template <typename T, int32_t D>
622 auto ArrayInterfaceStr(TensorView<T const, D> const &t) {
623 std::string str;
624 Json::Dump(ArrayInterface(t), &str);
625 return str;
626 }
627
628 template <typename T, int32_t D>
629 auto ArrayInterfaceStr(TensorView<T, D> const &t) {
630 std::string str;
631 Json::Dump(ArrayInterface(t), &str);
632 return str;
633 }
634
639 template <typename T, int32_t kDim = 5>
640 class Tensor {
641 public:
642 using ShapeT = size_t[kDim];
643 using StrideT = ShapeT;
644
645 private:
646 HostDeviceVector<T> data_;
647 ShapeT shape_{0};
648
649 template <typename I, std::int32_t D>
650 void Initialize(I const (&shape)[D], std::int32_t device) {
651 static_assert(D <= kDim, "Invalid shape.");
652 std::copy(shape, shape + D, shape_);
653 for (auto i = D; i < kDim; ++i) {
654 shape_[i] = 1;
655 }
656 if (device >= 0) {
657 data_.SetDevice(device);
658 data_.DevicePointer(); // Pull to device;
659 }
660 CHECK_EQ(data_.Size(), detail::CalcSize(shape_));
661 }
662
663 public:
664 Tensor() = default;
665
672 template <typename I, int32_t D>
673 explicit Tensor(I const (&shape)[D], int32_t device) {
674 // No device unroll as this is a host only function.
675 std::copy(shape, shape + D, shape_);
676 for (auto i = D; i < kDim; ++i) {
677 shape_[i] = 1;
678 }
679 auto size = detail::CalcSize(shape_);
680 if (device >= 0) {
681 data_.SetDevice(device);
682 }
683 data_.Resize(size);
684 if (device >= 0) {
685 data_.DevicePointer(); // Pull to device
686 }
687 }
691 template <typename It, typename I, int32_t D>
692 explicit Tensor(It begin, It end, I const (&shape)[D], int32_t device) {
693 auto &h_vec = data_.HostVector();
694 h_vec.insert(h_vec.begin(), begin, end);
695 // shape
696 this->Initialize(shape, device);
697 }
698
699 template <typename I, int32_t D>
700 explicit Tensor(std::initializer_list<T> data, I const (&shape)[D], int32_t device) {
701 auto &h_vec = data_.HostVector();
702 h_vec = data;
703 // shape
704 this->Initialize(shape, device);
705 }
706
710 TensorView<T, kDim> View(int32_t device) {
711 if (device >= 0) {
712 data_.SetDevice(device);
713 auto span = data_.DeviceSpan();
714 return {span, shape_, device};
715 } else {
716 auto span = data_.HostSpan();
717 return {span, shape_, device};
718 }
719 }
720 TensorView<T const, kDim> View(int32_t device) const {
721 if (device >= 0) {
722 data_.SetDevice(device);
723 auto span = data_.ConstDeviceSpan();
724 return {span, shape_, device};
725 } else {
726 auto span = data_.ConstHostSpan();
727 return {span, shape_, device};
728 }
729 }
730
731 auto HostView() const { return this->View(-1); }
732 auto HostView() { return this->View(-1); }
733
734 size_t Size() const { return data_.Size(); }
735 auto Shape() const { return common::Span<size_t const, kDim>{shape_}; }
736 auto Shape(size_t i) const { return shape_[i]; }
737
738 HostDeviceVector<T> *Data() { return &data_; }
739 HostDeviceVector<T> const *Data() const { return &data_; }
740
747 template <typename Fn>
748 void ModifyInplace(Fn &&fn) {
749 fn(this->Data(), common::Span<size_t, kDim>{this->shape_});
750 CHECK_EQ(this->Data()->Size(), detail::CalcSize(this->shape_))
751 << "Inconsistent size after modification.";
752 }
753
759 template <typename... S>
760 void Reshape(S &&...s) {
761 static_assert(sizeof...(S) <= kDim, "Invalid shape.");
762 detail::ReshapeImpl<0>(shape_, std::forward<S>(s)...);
763 auto constexpr kEnd = sizeof...(S);
764 static_assert(kEnd <= kDim, "Invalid shape.");
765 std::fill(shape_ + kEnd, shape_ + kDim, 1);
766 auto n = detail::CalcSize(shape_);
767 data_.Resize(n);
768 }
769
775 template <int32_t D>
776 void Reshape(size_t (&shape)[D]) {
777 static_assert(D <= kDim, "Invalid shape.");
778 std::copy(shape, shape + D, this->shape_);
779 std::fill(shape_ + D, shape_ + kDim, 1);
780 auto n = detail::CalcSize(shape_);
781 data_.Resize(n);
782 }
783
787 void SetDevice(int32_t device) const { data_.SetDevice(device); }
788 int32_t DeviceIdx() const { return data_.DeviceIdx(); }
789 };
790
791 // Only first axis is supported for now.
792 template <typename T, int32_t D>
793 void Stack(Tensor<T, D> *l, Tensor<T, D> const &r) {
794 if (r.DeviceIdx() >= 0) {
795 l->SetDevice(r.DeviceIdx());
796 }
797 l->ModifyInplace([&](HostDeviceVector<T> *data, common::Span<size_t, D> shape) {
798 for (size_t i = 1; i < D; ++i) {
799 if (shape[i] == 0) {
800 shape[i] = r.Shape(i);
801 } else {
802 CHECK_EQ(shape[i], r.Shape(i));
803 }
804 }
805 data->Extend(*r.Data());
806 shape[0] = l->Shape(0) + r.Shape(0);
807 });
808 }
809 } // namespace linalg
810 } // namespace xgboost
811
812 #if defined(LINALG_HD)
813 #undef LINALG_HD
814 #endif // defined(LINALG_HD)
815 #endif // XGBOOST_LINALG_H_
xgboost::linalg::detail::RangeTag::beg
I beg
Definition: linalg.h:79
xgboost::linalg::TensorView
A tensor view with static type and dimension. It implements indexing and slicing.
Definition: linalg.h:244
xgboost::JsonBoolean
Describes both true and false.
Definition: json.h:307
xgboost::linalg::TensorView::Size
LINALG_HD size_t Size() const
Number of items in the tensor.
Definition: linalg.h:467
xgboost::linalg::Tensor::Tensor
Tensor()=default
xgboost::HostDeviceVector::Extend
void Extend(const HostDeviceVector< T > &other)
xgboost::linalg::Tensor::ModifyInplace
void ModifyInplace(Fn &&fn)
Visitor function for modification that changes shape and data.
Definition: linalg.h:748
xgboost::linalg::TensorView::FContiguous
LINALG_HD bool FContiguous() const
Whether it's a f-contiguous array.
Definition: linalg.h:487
xgboost::linalg::Tensor< float, 2 >::ShapeT
size_t[kDim] ShapeT
Definition: linalg.h:642
xgboost::linalg::detail::Popc
LINALG_HD int Popc(uint32_t v)
Definition: linalg.h:129
xgboost::linalg::detail::ReshapeImpl
void ReshapeImpl(size_t(&out_shape)[D], I s)
Definition: linalg.h:187
xgboost::linalg::TensorView::Shape
LINALG_HD auto Shape(size_t i) const
Definition: linalg.h:457
xgboost::linalg::detail::Offset
constexpr size_t Offset(S(&strides)[D], size_t n, Head head)
Definition: linalg.h:46
xgboost::linalg::TensorView::ShapeT
size_t[kDim] ShapeT
Definition: linalg.h:246
xgboost::HostDeviceVector
Definition: host_device_vector.h:86
host_device_vector.h
A device-and-host vector abstraction layer.
xgboost::HostDeviceVector::DeviceSpan
common::Span< T > DeviceSpan()
xgboost::linalg::TensorView::Values
LINALG_HD auto Values() const -> decltype(data_) const &
Obtain a reference to the raw data.
Definition: linalg.h:497
xgboost::linalg::All
constexpr detail::AllTag All()
Specify all elements in the axis for slicing.
Definition: linalg.h:221
xgboost::linalg::TensorView::operator()
LINALG_HD T & operator()(Index &&...index)
Index the tensor to obtain a scalar value.
Definition: linalg.h:409
xgboost::linalg::detail::CalcStride
constexpr void CalcStride(size_t const (&shape)[D], size_t(&stride)[D])
Definition: linalg.h:59
xgboost::linalg::Tensor::DeviceIdx
int32_t DeviceIdx() const
Definition: linalg.h:788
xgboost::linalg::Tensor::Tensor
Tensor(std::initializer_list< T > data, I const (&shape)[D], int32_t device)
Definition: linalg.h:700
xgboost::linalg::ArrayInterface
Json ArrayInterface(TensorView< T const, D > const &t)
Array Interface defined by numpy.
Definition: linalg.h:578
xgboost::HostDeviceVector::DevicePointer
T * DevicePointer()
base.h
defines configuration macros of xgboost.
xgboost::HostDeviceVector::ConstHostSpan
common::Span< T const > ConstHostSpan() const
Definition: host_device_vector.h:114
xgboost::linalg::ArrayInterfaceStr
auto ArrayInterfaceStr(TensorView< T const, D > const &t)
Return string representation of array interface.
Definition: linalg.h:622
xgboost::linalg::Tensor::Data
const HostDeviceVector< T > * Data() const
Definition: linalg.h:739
xgboost::linalg::detail::UnrollLoop
constexpr LINALG_HD auto UnrollLoop(Fn fn)
Definition: linalg.h:113
xgboost::Json::Dump
static void Dump(Json json, std::string *out, std::ios::openmode mode=std::ios::out)
Encode the JSON object. Optional parameter mode for choosing between text and binary (ubjson) output.
xgboost::linalg::detail::AllTag
Definition: linalg.h:73
xgboost::linalg::TensorView::CContiguous
LINALG_HD bool CContiguous() const
Whether it's a c-contiguous array.
Definition: linalg.h:477
xgboost::linalg::Tensor::Shape
auto Shape(size_t i) const
Definition: linalg.h:736
xgboost::linalg::Tensor::Data
HostDeviceVector< T > * Data()
Definition: linalg.h:738
xgboost::linalg::detail::RangeTag::end
I end
Definition: linalg.h:80
xgboost::common::Span::empty
constexpr XGBOOST_DEVICE bool empty() const __span_noexcept
Definition: span.h:560
xgboost::HostDeviceVector::HostVector
std::vector< T > & HostVector()
xgboost::HostDeviceVector::SetDevice
void SetDevice(int device) const
xgboost::linalg::TensorView::kValueSize
constexpr static size_t kValueSize
Definition: linalg.h:338
xgboost::linalg::MakeTensorView
auto MakeTensorView(Container &data, I const (&shape)[D], int32_t device)
Constructor for automatic type deduction.
Definition: linalg.h:509
LINALG_HD
#define LINALG_HD
Definition: linalg.h:28
xgboost::JsonNull
Definition: json.h:291
xgboost::linalg::Tensor::SetDevice
void SetDevice(int32_t device) const
Set device ordinal for this tensor.
Definition: linalg.h:787
xgboost::linalg::detail::NativePopc
int32_t NativePopc(T v)
Definition: linalg.h:123
span.h
xgboost::linalg::TensorView::Slice
LINALG_HD auto Slice(S &&...slices) const
Slice the tensor. The returned tensor has inferred dim and shape. Scalar result is not supported.
Definition: linalg.h:440
xgboost::linalg::Stack
void Stack(Tensor< T, D > *l, Tensor< T, D > const &r)
Definition: linalg.h:793
xgboost::linalg::Tensor::View
TensorView< T const, kDim > View(int32_t device) const
Definition: linalg.h:720
xgboost::linalg::TensorView::TensorView
LINALG_HD TensorView(common::Span< T > data, I const (&shape)[D], I const (&stride)[D], int32_t device)
Create a tensor with data, shape and strides. Don't use this constructor if stride can be calculated ...
Definition: linalg.h:373
xgboost::linalg::detail::UnravelImpl
LINALG_HD auto UnravelImpl(I idx, common::Span< size_t const, D > shape)
Definition: linalg.h:167
xgboost::linalg::Tensor< float, 2 >::StrideT
ShapeT StrideT
Definition: linalg.h:643
xgboost::linalg::MakeVec
auto MakeVec(T *ptr, size_t s, int32_t device=-1)
Create a vector view from contigious memory.
Definition: linalg.h:547
xgboost::linalg::Tensor::View
TensorView< T, kDim > View(int32_t device)
Get a TensorView for this tensor.
Definition: linalg.h:710
xgboost::HostDeviceVector::HostSpan
common::Span< T > HostSpan()
Definition: host_device_vector.h:112
xgboost::linalg::Tensor::Reshape
void Reshape(S &&...s)
Reshape the tensor.
Definition: linalg.h:760
xgboost::common::Span::subspan
XGBOOST_DEVICE auto subspan() const -> Span< element_type, detail::ExtentValue< Extent, Offset, Count >::value >
Definition: span.h:595
xgboost::linalg::detail::IntTag
Definition: linalg.h:75
xgboost::linalg::detail::ArrayInterfaceHandler
Definition: linalg.h:36
xgboost::linalg::detail::Apply
decltype(auto) constexpr LINALG_HD Apply(Fn &&f, Tup &&t, std::index_sequence< I... >)
Definition: linalg.h:201
xgboost::linalg::Tensor
A tensor storage. To use it for other functionality like slicing one needs to obtain a view first....
Definition: linalg.h:640
xgboost::linalg::detail::RemoveCRType
std::remove_const_t< std::remove_reference_t< S > > RemoveCRType
Definition: linalg.h:107
xgboost::linalg::TensorView::operator()
LINALG_HD const T & operator()(Index &&...index) const
Index the tensor to obtain a scalar value.
Definition: linalg.h:419
xgboost::linalg::detail::ArrayInterfaceHandler::TypeChar
static constexpr char TypeChar()
Definition: linalg.h:38
xgboost::HostDeviceVector::Size
size_t Size() const
xgboost::linalg::Tensor::Shape
auto Shape() const
Definition: linalg.h:735
xgboost::linalg::TensorView::Contiguous
LINALG_HD bool Contiguous() const
Whether this is a contiguous array, both C and F contiguous returns true.
Definition: linalg.h:471
xgboost::HostDeviceVector::HostPointer
T * HostPointer()
Definition: host_device_vector.h:111
xgboost::HostDeviceVector::Resize
void Resize(size_t new_size, T v=T())
xgboost::linalg::UnravelIndex
LINALG_HD auto UnravelIndex(size_t idx, common::Span< size_t const, D > shape)
Turns linear index into multi-dimension index. Similar to numpy unravel.
Definition: linalg.h:523
xgboost::linalg::Tensor::HostView
auto HostView() const
Definition: linalg.h:731
xgboost::linalg::TensorView::DeviceIdx
LINALG_HD auto DeviceIdx() const
Obtain the CUDA device ordinal.
Definition: linalg.h:501
xgboost::JsonInteger
Definition: json.h:248
xgboost::HostDeviceVector::ConstDevicePointer
const T * ConstDevicePointer() const
xgboost::linalg::Tensor::HostView
auto HostView()
Definition: linalg.h:732
xgboost::linalg::Tensor::Reshape
void Reshape(size_t(&shape)[D])
Reshape the tensor.
Definition: linalg.h:776
xgboost::linalg::TensorView::kDimension
constexpr static size_t kDimension
Definition: linalg.h:339
xgboost::common::Span
span class implementation, based on ISO++20 span<T>. The interface should be the same.
Definition: span.h:148
xgboost::linalg::TensorView::TensorView
LINALG_HD TensorView(TensorView< U, kDim > const &that)
Definition: linalg.h:387
std
Definition: intrusive_ptr.h:207
xgboost::HostDeviceVector::ConstHostPointer
const T * ConstHostPointer() const
Definition: host_device_vector.h:115
xgboost::Integer
JsonInteger Integer
Definition: json.h:591
json.h
xgboost::common::Span::size
constexpr XGBOOST_DEVICE index_type size() const __span_noexcept
Definition: span.h:553
xgboost::JsonObject
Definition: json.h:187
xgboost::HostDeviceVector::DeviceIdx
int DeviceIdx() const
xgboost::JsonArray
Definition: json.h:110
xgboost::linalg::detail::CalcSize
constexpr size_t CalcSize(size_t(&shape)[D])
Definition: linalg.h:98
xgboost::common::Span::data
constexpr XGBOOST_DEVICE pointer data() const __span_noexcept
Definition: span.h:548
xgboost::linalg::TensorView::Stride
LINALG_HD auto Stride() const
Definition: linalg.h:458
xgboost::linalg::detail::RangeTag::Size
constexpr size_t Size() const
Definition: linalg.h:81
xgboost::JsonString
Definition: json.h:84
xgboost::linalg::Range
constexpr detail::RangeTag< I > Range(I beg, I end)
Specify a range of elements in the axis for slicing.
Definition: linalg.h:226
xgboost::linalg::Tensor::Tensor
Tensor(It begin, It end, I const (&shape)[D], int32_t device)
Definition: linalg.h:692
xgboost::linalg::detail::CalcSliceDim
constexpr int32_t CalcSliceDim()
Calculate the dimension of sliced tensor.
Definition: linalg.h:88
xgboost::HostDeviceVector::ConstDeviceSpan
common::Span< const T > ConstDeviceSpan() const
xgboost::Json
Data structure representing JSON format.
Definition: json.h:352
xgboost::linalg::detail::IndexToTag
std::conditional_t< std::is_integral< RemoveCRType< S > >::value, IntTag, S > IndexToTag
Definition: linalg.h:110
xgboost::linalg::Tensor::Size
size_t Size() const
Definition: linalg.h:734
xgboost::linalg::TensorView::TensorView
LINALG_HD TensorView(common::Span< T > data, I const (&shape)[D], int32_t device)
Create a tensor with data and shape.
Definition: linalg.h:354
xgboost::linalg::TensorView::Stride
LINALG_HD auto Stride(size_t i) const
Definition: linalg.h:462
xgboost::linalg::detail::Arr2Tup
constexpr auto Arr2Tup(T(&arr)[N], std::index_sequence< Idx... >)
Definition: linalg.h:154
xgboost::linalg::detail::RangeTag
Definition: linalg.h:78
xgboost::linalg::Tensor::Tensor
Tensor(I const (&shape)[D], int32_t device)
Create a tensor with shape and device ordinal. The storage is initialized automatically.
Definition: linalg.h:673
xgboost::linalg::TensorView::Shape
LINALG_HD auto Shape() const
Definition: linalg.h:453
xgboost
namespace of xgboost
Definition: base.h:110
xgboost::linalg::TensorView::StrideT
ShapeT StrideT
Definition: linalg.h:247