/
span.h
671 lines (569 loc) · 22.5 KB
/
span.h
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
/*!
* Copyright 2018 XGBoost contributors
* \brief span class based on ISO++20 span
*
* About NOLINTs in this file:
*
* If we want Span to work with std interface, like range for loop, the
* naming must be consistant with std, not XGBoost. Also, the interface also
* conflicts with XGBoost coding style, specifically, the use of `explicit'
* keyword.
*
*
* Some of the code is copied from Guidelines Support Library, here is the
* license:
*
* Copyright (c) 2015 Microsoft Corporation. All rights reserved.
*
* This code is licensed under the MIT License (MIT).
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
* THE SOFTWARE.
*/
#ifndef XGBOOST_SPAN_H_
#define XGBOOST_SPAN_H_
#include <xgboost/base.h>
#include <xgboost/logging.h>
#include <cinttypes> // size_t
#include <limits> // numeric_limits
#include <iterator>
#include <type_traits>
#include <cstdio>
/*!
* The version number 1910 is picked up from GSL.
*
* We might want to use MOODYCAMEL_NOEXCEPT from dmlc/concurrentqueue.h. But
* there are a lot more definitions in that file would cause warnings/troubles
* in MSVC 2013. Currently we try to keep the closure of Span as minimal as
* possible.
*
* There are other workarounds for MSVC, like _Unwrapped, _Verify_range ...
* Some of these are hiden magics of MSVC and I tried to avoid them. Should any
* of them become needed, please consult the source code of GSL, and possibily
* some explanations from this thread:
*
* https://github.com/Microsoft/GSL/pull/664
*
* TODO(trivialfis): Group these MSVC workarounds into a manageable place.
*/
#if defined(_MSC_VER) && _MSC_VER < 1910
#define __span_noexcept
#pragma push_macro("constexpr")
#define constexpr /*constexpr*/
#else
#define __span_noexcept noexcept
#endif // defined(_MSC_VER) && _MSC_VER < 1910
namespace xgboost {
namespace common {
// Usual logging facility is not available inside device code.
// assert is not supported in mac as of CUDA 10.0
#define KERNEL_CHECK(cond) \
do { \
if (!(cond)) { \
printf("\nKernel error:\n" \
"In: %s: %d\n" \
"\t%s\n\tExpecting: %s\n" \
"\tBlock: [%d, %d, %d], Thread: [%d, %d, %d]\n\n", \
__FILE__, __LINE__, __PRETTY_FUNCTION__, #cond, blockIdx.x, \
blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z); \
asm("trap;"); \
} \
} while (0);
#if defined(__CUDA_ARCH__)
#define SPAN_CHECK KERNEL_CHECK
#elif defined(XGBOOST_STRICT_R_MODE) && XGBOOST_STRICT_R_MODE == 1 // R package
#define SPAN_CHECK CHECK // check from dmlc
#else // not CUDA, not R
#define SPAN_CHECK(cond) \
do { \
if (XGBOOST_EXPECT(!(cond), false)) { \
fprintf(stderr, "[xgboost] Condition %s failed.\n", #cond); \
fflush(stderr); /* It seems stderr on Windows is beffered? */ \
std::terminate(); \
} \
} while (0);
#endif // __CUDA_ARCH__
#if defined(__CUDA_ARCH__)
#define SPAN_LT(lhs, rhs) \
if (!((lhs) < (rhs))) { \
printf("%lu < %lu failed\n", static_cast<size_t>(lhs), \
static_cast<size_t>(rhs)); \
asm("trap;"); \
}
#else
#define SPAN_LT(lhs, rhs) \
SPAN_CHECK((lhs) < (rhs))
#endif // defined(__CUDA_ARCH__)
namespace detail {
/*!
* By default, XGBoost uses uint32_t for indexing data. int64_t covers all
* values uint32_t can represent. Also, On x86-64 Linux, GCC uses long int to
* represent ptrdiff_t, which is just int64_t. So we make it determinstic
* here.
*/
using ptrdiff_t = typename std::conditional< // NOLINT
std::is_same<std::ptrdiff_t, std::int64_t>::value,
std::ptrdiff_t, std::int64_t>::type;
} // namespace detail
#if defined(_MSC_VER) && _MSC_VER < 1910
constexpr const std::size_t
dynamic_extent = std::numeric_limits<std::size_t>::max(); // NOLINT
#else
constexpr std::size_t dynamic_extent = std::numeric_limits<std::size_t>::max(); // NOLINT
#endif // defined(_MSC_VER) && _MSC_VER < 1910
enum class byte : unsigned char {}; // NOLINT
template <class ElementType, std::size_t Extent>
class Span;
namespace detail {
template <typename SpanType, bool IsConst>
class SpanIterator {
using ElementType = typename SpanType::element_type;
public:
using iterator_category = std::random_access_iterator_tag; // NOLINT
using value_type = typename SpanType::value_type; // NOLINT
using difference_type = detail::ptrdiff_t; // NOLINT
using reference = typename std::conditional< // NOLINT
IsConst, const ElementType, ElementType>::type&;
using pointer = typename std::add_pointer<reference>::type; // NOLINT
constexpr SpanIterator() = default;
XGBOOST_DEVICE constexpr SpanIterator(
const SpanType* _span,
typename SpanType::index_type _idx) __span_noexcept :
span_(_span), index_(_idx) {}
friend SpanIterator<SpanType, true>;
template <bool B, typename std::enable_if<!B && IsConst>::type* = nullptr>
XGBOOST_DEVICE constexpr SpanIterator( // NOLINT
const SpanIterator<SpanType, B>& other_) __span_noexcept
: SpanIterator(other_.span_, other_.index_) {}
XGBOOST_DEVICE reference operator*() const {
SPAN_CHECK(index_ < span_->size());
return *(span_->data() + index_);
}
XGBOOST_DEVICE reference operator[](difference_type n) const {
return *(*this + n);
}
XGBOOST_DEVICE pointer operator->() const {
SPAN_CHECK(index_ != span_->size());
return span_->data() + index_;
}
XGBOOST_DEVICE SpanIterator& operator++() {
SPAN_CHECK(index_ != span_->size());
index_++;
return *this;
}
XGBOOST_DEVICE SpanIterator operator++(int) {
auto ret = *this;
++(*this);
return ret;
}
XGBOOST_DEVICE SpanIterator& operator--() {
SPAN_CHECK(index_ != 0 && index_ <= span_->size());
index_--;
return *this;
}
XGBOOST_DEVICE SpanIterator operator--(int) {
auto ret = *this;
--(*this);
return ret;
}
XGBOOST_DEVICE SpanIterator operator+(difference_type n) const {
auto ret = *this;
return ret += n;
}
XGBOOST_DEVICE SpanIterator& operator+=(difference_type n) {
SPAN_CHECK((index_ + n) <= span_->size());
index_ += n;
return *this;
}
XGBOOST_DEVICE difference_type operator-(SpanIterator rhs) const {
SPAN_CHECK(span_ == rhs.span_);
return index_ - rhs.index_;
}
XGBOOST_DEVICE SpanIterator operator-(difference_type n) const {
auto ret = *this;
return ret -= n;
}
XGBOOST_DEVICE SpanIterator& operator-=(difference_type n) {
return *this += -n;
}
// friends
XGBOOST_DEVICE constexpr friend bool operator==(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return _lhs.span_ == _rhs.span_ && _lhs.index_ == _rhs.index_;
}
XGBOOST_DEVICE constexpr friend bool operator!=(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return !(_lhs == _rhs);
}
XGBOOST_DEVICE constexpr friend bool operator<(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return _lhs.index_ < _rhs.index_;
}
XGBOOST_DEVICE constexpr friend bool operator<=(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return !(_rhs < _lhs);
}
XGBOOST_DEVICE constexpr friend bool operator>(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return _rhs < _lhs;
}
XGBOOST_DEVICE constexpr friend bool operator>=(
SpanIterator _lhs, SpanIterator _rhs) __span_noexcept {
return !(_rhs > _lhs);
}
protected:
const SpanType *span_ { nullptr };
typename SpanType::index_type index_ { 0 };
};
// It's tempting to use constexpr instead of structs to do the following meta
// programming. But remember that we are supporting MSVC 2013 here.
/*!
* The extent E of the span returned by subspan is determined as follows:
*
* - If Count is not dynamic_extent, Count;
* - Otherwise, if Extent is not dynamic_extent, Extent - Offset;
* - Otherwise, dynamic_extent.
*/
template <std::size_t Extent, std::size_t Offset, std::size_t Count>
struct ExtentValue : public std::integral_constant<
std::size_t, Count != dynamic_extent ?
Count : (Extent != dynamic_extent ? Extent - Offset : Extent)> {};
/*!
* If N is dynamic_extent, the extent of the returned span E is also
* dynamic_extent; otherwise it is std::size_t(sizeof(T)) * N.
*/
template <typename T, std::size_t Extent>
struct ExtentAsBytesValue : public std::integral_constant<
std::size_t,
Extent == dynamic_extent ?
Extent : sizeof(T) * Extent> {};
template <std::size_t From, std::size_t To>
struct IsAllowedExtentConversion : public std::integral_constant<
bool, From == To || From == dynamic_extent || To == dynamic_extent> {};
template <class From, class To>
struct IsAllowedElementTypeConversion : public std::integral_constant<
bool, std::is_convertible<From(*)[], To(*)[]>::value> {};
template <class T>
struct IsSpanOracle : std::false_type {};
template <class T, std::size_t Extent>
struct IsSpanOracle<Span<T, Extent>> : std::true_type {};
template <class T>
struct IsSpan : public IsSpanOracle<typename std::remove_cv<T>::type> {};
// Re-implement std algorithms here to adopt CUDA.
template <typename T>
struct Less {
XGBOOST_DEVICE constexpr bool operator()(const T& _x, const T& _y) const {
return _x < _y;
}
};
template <typename T>
struct Greater {
XGBOOST_DEVICE constexpr bool operator()(const T& _x, const T& _y) const {
return _x > _y;
}
};
template <class InputIt1, class InputIt2,
class Compare =
detail::Less<decltype(std::declval<InputIt1>().operator*())>>
XGBOOST_DEVICE bool LexicographicalCompare(InputIt1 first1, InputIt1 last1,
InputIt2 first2, InputIt2 last2) {
Compare comp;
for (; first1 != last1 && first2 != last2; ++first1, ++first2) {
if (comp(*first1, *first2)) {
return true;
}
if (comp(*first2, *first1)) {
return false;
}
}
return first1 == last1 && first2 != last2;
}
} // namespace detail
/*!
* \brief span class implementation, based on ISO++20 span<T>. The interface
* should be the same.
*
* What's different from span<T> in Guidelines Support Library (GSL)
*
* Interface might be slightly different, we stick with ISO.
*
* GSL uses C++14/17 features, which are not available here.
* GSL uses constexpr extensively, which is not possibile with limitation
* of C++11.
* GSL doesn't concern about CUDA.
*
* GSL is more thoroughly implemented and tested.
* GSL is more optimized, especially for static extent.
*
* GSL uses __buildin_unreachable() when error, Span<T> uses dmlc LOG and
* customized CUDA logging.
*
*
* What's different from span<T> in ISO++20 (ISO)
*
* ISO uses functions/structs from std library, which might be not available
* in CUDA.
* Initializing from std::array is not supported.
*
* ISO uses constexpr extensively, which is not possibile with limitation
* of C++11.
* ISO uses C++14/17 features, which is not available here.
* ISO doesn't concern about CUDA.
*
* ISO uses std::terminate(), Span<T> uses dmlc LOG and customized CUDA
* logging.
*
*
* Limitations:
* With thrust:
* It's not adviced to initialize Span with host_vector directly, since
* host_vector::data() is a host function.
* It's not possible to initialize Span with device_vector directly, since
* device_vector::data() returns a wrapped pointer.
* It's unclear that what kind of thrust algorithm can be used without
* memory error. See the test case "GPUSpan.WithTrust"
*
* Pass iterator to kernel:
* Not possible. Use subspan instead.
*
* The underlying Span in SpanIterator is a pointer, but CUDA pass kernel
* parameter by value. If we were to hold a Span value instead of a
* pointer, the following snippet will crash, violating the safety
* purpose of Span:
*
* \code{.cpp}
* Span<float> span {arr_a};
* auto beg = span.begin();
*
* Span<float> span_b = arr_b;
* span = span_b;
*
* delete arr_a;
* beg++; // crash
* \endcode
*
* While hoding a pointer or reference should avoid the problem, its a
* compromise. Since we have subspan, it's acceptable not to support
* passing iterator.
*/
template <typename T,
std::size_t Extent = dynamic_extent>
class Span {
public:
using element_type = T; // NOLINT
using value_type = typename std::remove_cv<T>::type; // NOLINT
using index_type = std::size_t; // NOLINT
using difference_type = detail::ptrdiff_t; // NOLINT
using pointer = T*; // NOLINT
using reference = T&; // NOLINT
using iterator = detail::SpanIterator<Span<T, Extent>, false>; // NOLINT
using const_iterator = const detail::SpanIterator<Span<T, Extent>, true>; // NOLINT
using reverse_iterator = detail::SpanIterator<Span<T, Extent>, false>; // NOLINT
using const_reverse_iterator = const detail::SpanIterator<Span<T, Extent>, true>; // NOLINT
// constructors
constexpr Span() __span_noexcept = default;
XGBOOST_DEVICE Span(pointer _ptr, index_type _count) :
size_(_count), data_(_ptr) {
SPAN_CHECK(!(Extent != dynamic_extent && _count != Extent));
SPAN_CHECK(_ptr || _count == 0);
}
XGBOOST_DEVICE Span(pointer _first, pointer _last) :
size_(_last - _first), data_(_first) {
SPAN_CHECK(data_ || size_ == 0);
}
template <std::size_t N>
XGBOOST_DEVICE constexpr Span(element_type (&arr)[N]) // NOLINT
__span_noexcept : size_(N), data_(&arr[0]) {}
template <class Container,
class = typename std::enable_if<
!std::is_const<element_type>::value &&
!detail::IsSpan<Container>::value &&
std::is_convertible<typename Container::pointer, pointer>::value &&
std::is_convertible<typename Container::pointer,
decltype(std::declval<Container>().data())>::value>::type>
Span(Container& _cont) : // NOLINT
size_(_cont.size()), data_(_cont.data()) {
static_assert(!detail::IsSpan<Container>::value, "Wrong constructor of Span is called.");
}
template <class Container,
class = typename std::enable_if<
std::is_const<element_type>::value &&
!detail::IsSpan<Container>::value &&
std::is_convertible<typename Container::pointer, pointer>::value &&
std::is_convertible<typename Container::pointer,
decltype(std::declval<Container>().data())>::value>::type>
Span(const Container& _cont) : size_(_cont.size()), // NOLINT
data_(_cont.data()) {
static_assert(!detail::IsSpan<Container>::value, "Wrong constructor of Span is called.");
}
template <class U, std::size_t OtherExtent,
class = typename std::enable_if<
detail::IsAllowedElementTypeConversion<U, T>::value &&
detail::IsAllowedExtentConversion<OtherExtent, Extent>::value>>
XGBOOST_DEVICE constexpr Span(const Span<U, OtherExtent>& _other) // NOLINT
__span_noexcept : size_(_other.size()), data_(_other.data()) {}
XGBOOST_DEVICE constexpr Span(const Span& _other)
__span_noexcept : size_(_other.size()), data_(_other.data()) {}
XGBOOST_DEVICE Span& operator=(const Span& _other) __span_noexcept {
size_ = _other.size();
data_ = _other.data();
return *this;
}
XGBOOST_DEVICE ~Span() __span_noexcept {}; // NOLINT
XGBOOST_DEVICE constexpr iterator begin() const __span_noexcept { // NOLINT
return {this, 0};
}
XGBOOST_DEVICE constexpr iterator end() const __span_noexcept { // NOLINT
return {this, size()};
}
XGBOOST_DEVICE constexpr const_iterator cbegin() const __span_noexcept { // NOLINT
return {this, 0};
}
XGBOOST_DEVICE constexpr const_iterator cend() const __span_noexcept { // NOLINT
return {this, size()};
}
XGBOOST_DEVICE constexpr reverse_iterator rbegin() const __span_noexcept { // NOLINT
return reverse_iterator{end()};
}
XGBOOST_DEVICE constexpr reverse_iterator rend() const __span_noexcept { // NOLINT
return reverse_iterator{begin()};
}
XGBOOST_DEVICE constexpr const_reverse_iterator crbegin() const __span_noexcept { // NOLINT
return const_reverse_iterator{cend()};
}
XGBOOST_DEVICE constexpr const_reverse_iterator crend() const __span_noexcept { // NOLINT
return const_reverse_iterator{cbegin()};
}
// element access
XGBOOST_DEVICE reference front() const { // NOLINT
return (*this)[0];
}
XGBOOST_DEVICE reference back() const { // NOLINT
return (*this)[size() - 1];
}
XGBOOST_DEVICE reference operator[](index_type _idx) const {
SPAN_LT(_idx, size());
return data()[_idx];
}
XGBOOST_DEVICE reference operator()(index_type _idx) const {
return this->operator[](_idx);
}
XGBOOST_DEVICE constexpr pointer data() const __span_noexcept { // NOLINT
return data_;
}
// Observers
XGBOOST_DEVICE constexpr index_type size() const __span_noexcept { // NOLINT
return size_;
}
XGBOOST_DEVICE constexpr index_type size_bytes() const __span_noexcept { // NOLINT
return size() * sizeof(T);
}
XGBOOST_DEVICE constexpr bool empty() const __span_noexcept { // NOLINT
return size() == 0;
}
// Subviews
template <std::size_t Count>
XGBOOST_DEVICE Span<element_type, Count> first() const { // NOLINT
SPAN_CHECK(Count <= size());
return {data(), Count};
}
XGBOOST_DEVICE Span<element_type, dynamic_extent> first( // NOLINT
std::size_t _count) const {
SPAN_CHECK(_count <= size());
return {data(), _count};
}
template <std::size_t Count>
XGBOOST_DEVICE Span<element_type, Count> last() const { // NOLINT
SPAN_CHECK(Count <= size());
return {data() + size() - Count, Count};
}
XGBOOST_DEVICE Span<element_type, dynamic_extent> last( // NOLINT
std::size_t _count) const {
SPAN_CHECK(_count <= size());
return subspan(size() - _count, _count);
}
/*!
* If Count is std::dynamic_extent, r.size() == this->size() - Offset;
* Otherwise r.size() == Count.
*/
template <std::size_t Offset,
std::size_t Count = dynamic_extent>
XGBOOST_DEVICE auto subspan() const -> // NOLINT
Span<element_type,
detail::ExtentValue<Extent, Offset, Count>::value> {
SPAN_CHECK((Count == dynamic_extent) ?
(Offset <= size()) : (Offset + Count <= size()));
return {data() + Offset, Count == dynamic_extent ? size() - Offset : Count};
}
XGBOOST_DEVICE Span<element_type, dynamic_extent> subspan( // NOLINT
index_type _offset,
index_type _count = dynamic_extent) const {
SPAN_CHECK((_count == dynamic_extent) ?
(_offset <= size()) : (_offset + _count <= size()));
return {data() + _offset, _count ==
dynamic_extent ? size() - _offset : _count};
}
private:
index_type size_ { 0 };
pointer data_ { nullptr };
};
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE bool operator==(Span<T, X> l, Span<U, Y> r) {
if (l.size() != r.size()) {
return false;
}
for (auto l_beg = l.cbegin(), r_beg = r.cbegin(); l_beg != l.cend();
++l_beg, ++r_beg) {
if (*l_beg != *r_beg) {
return false;
}
}
return true;
}
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE constexpr bool operator!=(Span<T, X> l, Span<U, Y> r) {
return !(l == r);
}
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE constexpr bool operator<(Span<T, X> l, Span<U, Y> r) {
return detail::LexicographicalCompare(l.begin(), l.end(),
r.begin(), r.end());
}
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE constexpr bool operator<=(Span<T, X> l, Span<U, Y> r) {
return !(l > r);
}
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE constexpr bool operator>(Span<T, X> l, Span<U, Y> r) {
return detail::LexicographicalCompare<
typename Span<T, X>::iterator, typename Span<U, Y>::iterator,
detail::Greater<typename Span<T, X>::element_type>>(l.begin(), l.end(),
r.begin(), r.end());
}
template <class T, std::size_t X, class U, std::size_t Y>
XGBOOST_DEVICE constexpr bool operator>=(Span<T, X> l, Span<U, Y> r) {
return !(l < r);
}
template <class T, std::size_t E>
XGBOOST_DEVICE auto as_bytes(Span<T, E> s) __span_noexcept -> // NOLINT
Span<const byte, detail::ExtentAsBytesValue<T, E>::value> {
return {reinterpret_cast<const byte*>(s.data()), s.size_bytes()};
}
template <class T, std::size_t E>
XGBOOST_DEVICE auto as_writable_bytes(Span<T, E> s) __span_noexcept -> // NOLINT
Span<byte, detail::ExtentAsBytesValue<T, E>::value> {
return {reinterpret_cast<byte*>(s.data()), s.size_bytes()};
}
} // namespace common
} // namespace xgboost
#if defined(_MSC_VER) &&_MSC_VER < 1910
#undef constexpr
#pragma pop_macro("constexpr")
#undef __span_noexcept
#endif // _MSC_VER < 1910
#endif // XGBOOST_SPAN_H_