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>
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 
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>
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);
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);
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>
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>
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>
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>
570 
577 template <typename T, int32_t D>
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>
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>
623  std::string str;
624  Json::Dump(ArrayInterface(t), &str);
625  return str;
626 }
627 
628 template <typename T, int32_t D>
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  }
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