Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

platform.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without modification, are permitted
5 * provided that the following conditions are met:
6 * * Redistributions of source code must retain the above copyright notice, this list of
7 * conditions and the following disclaimer.
8 * * Redistributions in binary form must reproduce the above copyright notice, this list of
9 * conditions and the following disclaimer in the documentation and/or other materials
10 * provided with the distribution.
11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12 * to endorse or promote products derived from this software without specific prior written
13 * permission.
14 *
15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23 *
24 **************************************************************************************************/
25
26 #pragma once
27
94 //-----------------------------------------------------------------------------
95 // Dependencies
96 //-----------------------------------------------------------------------------
97
98 #include <stdint.h>
99
100 #if !defined(__CUDACC_RTC__)
101 //-----------------------------------------------------------------------------
102 // Include STL files that platform provides functionality for
103 //-----------------------------------------------------------------------------
104
105 #include <algorithm> // Minimum/maximum operations
106 #include <cstddef> // nullptr_t
107 #include <functional> // Arithmetic operations
108 #include <utility> // For methods on std::pair
109 #if (!defined(_MSC_VER) && (__cplusplus >= 201103L)) || (defined(_MSC_VER) && (_MS_VER >= 1500))
110 #include <type_traits> // For integral constants, conditional metaprogramming, and type traits
111 #endif
112
113 #include "cutlass/cutlass.h"
114
115 #endif
116
117 //-----------------------------------------------------------------------------
118 // OS
119 //-----------------------------------------------------------------------------
120 #if defined(WIN32) || defined(_WIN32) || defined(__WIN32) && !defined(__CYGWIN__)
121 #define CUTLASS_OS_WINDOWS
122 #endif
123
124 /******************************************************************************
125 * Macros
126 ******************************************************************************/
127 //-----------------------------------------------------------------------------
128 // Keywords
129 //-----------------------------------------------------------------------------
130
132 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
133 #ifndef noexcept
134 #define noexcept
135 #endif
136 #ifndef constexpr
137 #define constexpr
138 #endif
139 #endif
140
142 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1310))
143 #ifndef nullptr
144 #define nullptr 0
145 #endif
146 #endif
147
149 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600))
150 #ifndef static_assert
151 #define __platform_cat_(a, b) a##b
152 #define __platform_cat(a, b) __platform_cat_(a, b)
153 #define static_assert(__e, __m) typedef int __platform_cat(AsSeRt, __LINE__)[(__e) ? 1 : -1]
154 #endif
155 #endif
156
157 //-----------------------------------------------------------------------------
158 // Functions
159 //-----------------------------------------------------------------------------
160
162 #ifndef __NV_STD_MAX
163 #define __NV_STD_MAX(a, b) (((b) > (a)) ? (b) : (a))
164 #endif
165
167 #ifndef __NV_STD_MIN
168 #define __NV_STD_MIN(a, b) (((b) < (a)) ? (b) : (a))
169 #endif
170
171 /******************************************************************************
172 * Re-implementations
173 ******************************************************************************/
174 namespace cutlass {
175 namespace platform {
176
177 //-----------------------------------------------------------------------------
178 // Arithmetic operations, comparisons <functional>
179 //-----------------------------------------------------------------------------
180
182 template <typename T>
183 struct plus {
184 CUTLASS_HOST_DEVICE constexpr T operator()(const T& lhs, const T& rhs) const { return lhs + rhs; }
185 };
186
188 template <typename T>
189 struct less {
190 CUTLASS_HOST_DEVICE constexpr bool operator()(const T& lhs, const T& rhs) const {
191 return lhs < rhs;
192 }
193 };
194
196 template <typename T>
197 struct greater {
198 CUTLASS_HOST_DEVICE constexpr bool operator()(const T& lhs, const T& rhs) const {
199 return lhs > rhs;
200 }
201 };
202
203 //-----------------------------------------------------------------------------
204 // Minimum/maximum operations <algorithm>
205 //-----------------------------------------------------------------------------
206
208 template <typename T>
209 CUTLASS_HOST_DEVICE constexpr const T& min(const T& a, const T& b) {
210 return (b < a) ? b : a;
211 }
212
214 template <typename T>
215 CUTLASS_HOST_DEVICE constexpr const T& max(const T& a, const T& b) {
216 return (a < b) ? b : a;
217 }
218
219 #if !defined(__CUDACC_RTC__)
220 //-----------------------------------------------------------------------------
221 // Methods on std::pair
222 //-----------------------------------------------------------------------------
223
224 using std::pair;
225
226 template <class T1, class T2>
227 CUTLASS_HOST_DEVICE constexpr bool operator==(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
228 return (lhs.first == rhs.first) && (lhs.second == rhs.second);
229 }
230
231 template <class T1, class T2>
232 CUTLASS_HOST_DEVICE constexpr bool operator!=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
233 return (lhs.first != rhs.first) && (lhs.second != rhs.second);
234 }
235
236 template <class T1, class T2>
237 CUTLASS_HOST_DEVICE constexpr bool operator<(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
238 return (lhs.first < rhs.first) ? true : (rhs.first < lhs.first) ? false
239 : (lhs.second < rhs.second);
240 }
241
242 template <class T1, class T2>
243 CUTLASS_HOST_DEVICE constexpr bool operator<=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
244 return !(rhs < lhs);
245 }
246
247 template <class T1, class T2>
248 CUTLASS_HOST_DEVICE constexpr bool operator>(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
249 return (rhs < lhs);
250 }
251
252 template <class T1, class T2>
253 CUTLASS_HOST_DEVICE constexpr bool operator>=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
254 return !(lhs < rhs);
255 }
256
257 template <class T1, class T2>
258 CUTLASS_HOST_DEVICE std::pair<T1, T2> make_pair(T1 t, T2 u) {
259 std::pair<T1, T2> retval;
260 retval.first = t;
261 retval.second = u;
262 return retval;
263 }
264 #endif
265
266 } // namespace platform
267
268 /******************************************************************************
269 * Implementations of C++ 11/14/17/... STL features
270 ******************************************************************************/
271
272 namespace platform {
273
274 //-----------------------------------------------------------------------------
275 // Integral constant helper types <type_traits>
276 //-----------------------------------------------------------------------------
277
278 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
279
281 template <typename value_t, value_t V>
282 struct integral_constant;
283
285 template <typename value_t, value_t V>
286 struct integral_constant {
287 static const value_t value = V;
288
289 typedef value_t value_type;
290 typedef integral_constant<value_t, V> type;
291
292 CUTLASS_HOST_DEVICE operator value_type() const { return value; }
293
294 CUTLASS_HOST_DEVICE const value_type operator()() const { return value; }
295 };
296
297 #else
298
299 using std::integral_constant;
300 using std::pair;
301
302 #endif
303
305 typedef integral_constant<bool, true> true_type;
306
308 typedef integral_constant<bool, false> false_type;
309
310 #if (!defined(_MSC_VER) && (__cplusplus <= 201402L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
311
313 template <bool V>
314 struct bool_constant : platform::integral_constant<bool, V> {};
315
316 #else
317
318 using std::bool_constant;
319
320 #endif
321
322 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1700))
323
325 struct nullptr_t {};
326
327 #else
328
329 using std::nullptr_t;
330
331 #endif
332
333 //-----------------------------------------------------------------------------
334 // Conditional metaprogramming <type_traits>
335 //-----------------------------------------------------------------------------
336
337 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600))
338
340 template <bool C, typename T = void>
341 struct enable_if {
342 typedef T type;
343 };
344
346 template <typename T>
347 struct enable_if<false, T> {};
348
350 template <bool B, class T, class F>
351 struct conditional {
352 typedef T type;
353 };
354
356 template <class T, class F>
357 struct conditional<false, T, F> {
358 typedef F type;
359 };
360
361 #else
362
363 using std::enable_if;
364 using std::conditional;
365
366 #endif
367
368 //-----------------------------------------------------------------------------
369 // Const/volatility specifiers <type_traits>
370 //-----------------------------------------------------------------------------
371
372 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
373
375 template <typename T>
376 struct remove_const {
377 typedef T type;
378 };
379
381 template <typename T>
382 struct remove_const<const T> {
383 typedef T type;
384 };
385
387 template <typename T>
388 struct remove_volatile {
389 typedef T type;
390 };
391
393 template <typename T>
394 struct remove_volatile<volatile T> {
395 typedef T type;
396 };
397
399 template <typename T>
400 struct remove_cv {
401 typedef typename remove_volatile<typename remove_const<T>::type>::type type;
402 };
403
404 #else
405
406 using std::remove_const;
407 using std::remove_volatile;
408 using std::remove_cv;
409
410 #endif
411
412 //-----------------------------------------------------------------------------
413 // Type relationships <type_traits>
414 //-----------------------------------------------------------------------------
415
416 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
417
419 template <typename A, typename B>
420 struct is_same : false_type {};
421
423 template <typename A>
424 struct is_same<A, A> : true_type {};
425
427 template <typename BaseT, typename DerivedT>
428 struct is_base_of_helper {
429 typedef char (&yes)[1];
430 typedef char (&no)[2];
431
432 template <typename B, typename D>
433 struct dummy {
434 CUTLASS_HOST_DEVICE operator B*() const;
435 CUTLASS_HOST_DEVICE operator D*();
436 };
437
438 template <typename T>
439 CUTLASS_HOST_DEVICE static yes check(DerivedT*, T);
440
441 CUTLASS_HOST_DEVICE static no check(BaseT*, int);
442
443 static const bool value = sizeof(check(dummy<BaseT, DerivedT>(), int())) == sizeof(yes);
444 };
445
447 template <typename BaseT, typename DerivedT>
448 struct is_base_of
449 : integral_constant<bool,
450 (is_base_of_helper<typename remove_cv<BaseT>::type,
451 typename remove_cv<DerivedT>::type>::value) ||
452 (is_same<typename remove_cv<BaseT>::type,
453 typename remove_cv<DerivedT>::type>::value)> {};
454
455 #else
456
457 using std::is_same;
458 using std::is_base_of;
459
460 #endif
461
462 //-----------------------------------------------------------------------------
463 // Type properties <type_traits>
464 //-----------------------------------------------------------------------------
465
466 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
467
469 template <typename T>
470 struct is_volatile : false_type {};
471 template <typename T>
472 struct is_volatile<volatile T> : true_type {};
473
475 template <typename T>
476 struct is_pointer_helper : false_type {};
477
479 template <typename T>
480 struct is_pointer_helper<T*> : true_type {};
481
483 template <typename T>
484 struct is_pointer : is_pointer_helper<typename remove_cv<T>::type> {};
485
487 template <typename T>
488 struct is_void : is_same<void, typename remove_cv<T>::type> {};
489
491 template <typename T>
492 struct is_integral : false_type {};
493 template <>
494 struct is_integral<char> : true_type {};
495 template <>
496 struct is_integral<signed char> : true_type {};
497 template <>
498 struct is_integral<unsigned char> : true_type {};
499 template <>
500 struct is_integral<short> : true_type {};
501 template <>
502 struct is_integral<unsigned short> : true_type {};
503 template <>
504 struct is_integral<int> : true_type {};
505 template <>
506 struct is_integral<unsigned int> : true_type {};
507 template <>
508 struct is_integral<long> : true_type {};
509 template <>
510 struct is_integral<unsigned long> : true_type {};
511 template <>
512 struct is_integral<long long> : true_type {};
513 template <>
514 struct is_integral<unsigned long long> : true_type {};
515 template <typename T>
516 struct is_integral<volatile T> : is_integral<T> {};
517 template <typename T>
518 struct is_integral<const T> : is_integral<T> {};
519 template <typename T>
520 struct is_integral<const volatile T> : is_integral<T> {};
521
523 template <typename T>
524 struct is_floating_point
525 : integral_constant<bool,
526 (is_same<float, typename remove_cv<T>::type>::value ||
527 is_same<double, typename remove_cv<T>::type>::value)> {};
528
530 template <typename T>
531 struct is_arithmetic
532 : integral_constant<bool, (is_integral<T>::value || is_floating_point<T>::value)> {};
533
535 template <typename T>
536 struct is_fundamental
537 : integral_constant<bool,
538 (is_arithmetic<T>::value || is_void<T>::value ||
539 is_same<nullptr_t, typename remove_cv<T>::type>::value)> {};
540
541 #else
542
543 using std::is_volatile;
544 using std::is_pointer;
545 using std::is_void;
546 using std::is_integral;
547 using std::is_floating_point;
548 using std::is_arithmetic;
549 using std::is_fundamental;
550
551 #endif
552
553 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800)) || \
554 (defined(__GNUG__) && (__GNUC__ < 5))
555
566 template <typename T>
567 struct is_trivially_copyable
568 : integral_constant<bool, (is_fundamental<T>::value || is_pointer<T>::value)> {};
569
570 #else
571
572 using std::is_trivially_copyable;
573
574 #endif
575
576 //-----------------------------------------------------------------------------
577 // Alignment and layout utilities
578 //-----------------------------------------------------------------------------
579
580 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
581
583 template <typename value_t>
584 struct alignment_of {
585 struct pad {
586 value_t val;
587 char byte;
588 };
589
590 enum { value = sizeof(pad) - sizeof(value_t) };
591 };
592
593 #else
594
595 template <typename value_t>
596 struct alignment_of : std::alignment_of<value_t> {};
597
598 #endif
599
600 /* 16B specializations where 32-bit Win32 host compiler disagrees with device compiler */
601 template <>
602 struct alignment_of<int4> {
603 enum { value = 16 };
604 };
605 template <>
606 struct alignment_of<uint4> {
607 enum { value = 16 };
608 };
609 template <>
610 struct alignment_of<float4> {
611 enum { value = 16 };
612 };
613 template <>
614 struct alignment_of<long4> {
615 enum { value = 16 };
616 };
617 template <>
618 struct alignment_of<ulong4> {
619 enum { value = 16 };
620 };
621 template <>
622 struct alignment_of<longlong2> {
623 enum { value = 16 };
624 };
625 template <>
626 struct alignment_of<ulonglong2> {
627 enum { value = 16 };
628 };
629 template <>
630 struct alignment_of<double2> {
631 enum { value = 16 };
632 };
633 template <>
634 struct alignment_of<longlong4> {
635 enum { value = 16 };
636 };
637 template <>
638 struct alignment_of<ulonglong4> {
639 enum { value = 16 };
640 };
641 template <>
642 struct alignment_of<double4> {
643 enum { value = 16 };
644 };
645
646 // Specializations for volatile/const qualified types
647 template <typename value_t>
648 struct alignment_of<volatile value_t> : alignment_of<value_t> {};
649 template <typename value_t>
650 struct alignment_of<const value_t> : alignment_of<value_t> {};
651 template <typename value_t>
652 struct alignment_of<const volatile value_t> : alignment_of<value_t> {};
653
654 #if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800))
655
656 template <size_t Align>
657 struct aligned_chunk;
658 template <>
659 struct __align__(1) aligned_chunk<1> {
660 uint8_t buff;
661 };
662 template <>
663 struct __align__(2) aligned_chunk<2> {
664 uint16_t buff;
665 };
666 template <>
667 struct __align__(4) aligned_chunk<4> {
668 uint32_t buff;
669 };
670 template <>
671 struct __align__(8) aligned_chunk<8> {
672 uint32_t buff[2];
673 };
674 template <>
675 struct __align__(16) aligned_chunk<16> {
676 uint32_t buff[4];
677 };
678 template <>
679 struct __align__(32) aligned_chunk<32> {
680 uint32_t buff[8];
681 };
682 template <>
683 struct __align__(64) aligned_chunk<64> {
684 uint32_t buff[16];
685 };
686 template <>
687 struct __align__(128) aligned_chunk<128> {
688 uint32_t buff[32];
689 };
690 template <>
691 struct __align__(256) aligned_chunk<256> {
692 uint32_t buff[64];
693 };
694 template <>
695 struct __align__(512) aligned_chunk<512> {
696 uint32_t buff[128];
697 };
698 template <>
699 struct __align__(1024) aligned_chunk<1024> {
700 uint32_t buff[256];
701 };
702 template <>
703 struct __align__(2048) aligned_chunk<2048> {
704 uint32_t buff[512];
705 };
706 template <>
707 struct __align__(4096) aligned_chunk<4096> {
708 uint32_t buff[1024];
709 };
710
712 template <size_t Len, size_t Align>
713 struct aligned_storage {
714 typedef aligned_chunk<Align> type[Len / sizeof(aligned_chunk<Align>)];
715 };
716
717 #else
718
719 using std::aligned_storage;
720
721 #endif
722
723 #if !defined(__CUDACC_RTC__)
724 template <typename T>
726 struct default_delete {
727 void operator()(T* ptr) const { delete ptr; }
728 };
729
731 template <typename T>
732 struct default_delete<T[]> {
733 void operator()(T* ptr) const { delete[] ptr; }
734 };
735
737 template <class T, class Deleter = default_delete<T> >
738 class unique_ptr {
739 public:
740 typedef T* pointer;
741 typedef T element_type;
742 typedef Deleter deleter_type;
743
744 private:
746 pointer _ptr;
747
749 deleter_type _deleter;
750
751 public:
752 unique_ptr() : _ptr(nullptr) {}
753 unique_ptr(pointer p) : _ptr(p) {}
754
755 ~unique_ptr() {
756 if (_ptr) {
757 _deleter(_ptr);
758 }
759 }
761 pointer get() const noexcept { return _ptr; }
762
764 pointer release() noexcept {
765 pointer p(_ptr);
766 _ptr = nullptr;
767 return p;
768 }
769
771 void reset(pointer p = pointer()) noexcept {
772 pointer old_ptr = _ptr;
773 _ptr = p;
774 if (old_ptr != nullptr) {
775 get_deleter()(old_ptr);
776 }
777 }
778
780 void swap(unique_ptr& other) noexcept { std::swap(_ptr, other._ptr); }
781
783 Deleter& get_deleter() noexcept { return _deleter; }
784
786 Deleter const& get_deleter() const noexcept { return _deleter; }
787
789 operator bool() const noexcept { return _ptr != nullptr; }
790
792 T& operator*() const { return *_ptr; }
793
795 pointer operator->() const noexcept { return _ptr; }
796
798 T& operator[](size_t i) const { return _ptr[i]; }
799 };
800
802 template <typename T, typename Deleter>
803 void swap(unique_ptr<T, Deleter>& lhs, unique_ptr<T, Deleter>& rhs) noexcept {
804 lhs.swap(rhs);
805 }
806 #endif
807
808 }; // namespace platform
809 }; // namespace cutlass
cutlass::platform::integral_constant::value
static const value_t value
Definition: platform.h:287
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:215
cutlass
Definition: convert.h:33
constexpr
#define constexpr
Definition: platform.h:137
cutlass::platform::nullptr_t
std::nullptr_t
Definition: platform.h:325
cutlass::platform::swap
void swap(unique_ptr< T, Deleter > &lhs, unique_ptr< T, Deleter > &rhs) noexcept
Specializes the swap algorithm.
Definition: platform.h:803
cutlass::platform::is_pointer_helper
Helper for std::is_pointer (false specialization)
Definition: platform.h:476
cutlass::platform::unique_ptr::deleter_type
Deleter deleter_type
Definition: platform.h:742
cutlass::platform::remove_const::type
T type
Definition: platform.h:377
cutlass::platform::alignment_of::pad::val
value_t val
Definition: platform.h:586
cutlass::platform::conditional::type
T type
Definition: platform.h:352
cutlass::platform::unique_ptr::pointer
T * pointer
Definition: platform.h:740
cutlass::platform::less
std::less
Definition: platform.h:189
cutlass::platform::is_same
std::is_same (false specialization)
Definition: platform.h:420
cutlass::platform::is_pointer
std::is_pointer
Definition: platform.h:484
cutlass::platform::integral_constant::value_type
value_t value_type
Definition: platform.h:289
cutlass::platform::make_pair
CUTLASS_HOST_DEVICE std::pair< T1, T2 > make_pair(T1 t, T2 u)
Definition: platform.h:258
cutlass::platform::unique_ptr::unique_ptr
unique_ptr()
Definition: platform.h:752
cutlass::platform::operator==
CUTLASS_HOST_DEVICE bool operator==(complex< T > const &lhs, complex< T > const &rhs)
Equality operator.
Definition: complex.h:224
cutlass::platform::greater
std::greater
Definition: platform.h:197
cutlass::platform::is_void
std::is_void
Definition: platform.h:488
cutlass::platform::unique_ptr::operator->
pointer operator->() const noexcept
Returns a pointer to the managed object.
Definition: platform.h:795
cutlass::platform::unique_ptr::operator[]
T & operator[](size_t i) const
Array access to managed object.
Definition: platform.h:798
cutlass::platform::default_delete< T[]>::operator()
void operator()(T *ptr) const
Definition: platform.h:733
cutlass::platform::default_delete
Default deleter.
Definition: platform.h:726
cutlass::platform::alignment_of::value
Definition: platform.h:590
cutlass::platform::unique_ptr
std::unique_ptr
Definition: platform.h:738
cutlass::platform::alignment_of::pad
Definition: platform.h:585
cutlass::platform::is_floating_point
std::is_floating_point
Definition: platform.h:524
cutlass::platform::conditional< false, T, F >::type
F type
Definition: platform.h:358
cutlass::platform::operator>
CUTLASS_HOST_DEVICE bool operator>(Pair< T1, T2 > const &lhs, Pair< T1, T2 > const &rhs)
Lexical comparison.
Definition: pair.h:110
cutlass::platform::false_type
integral_constant< bool, false > false_type
The type used as a compile-time boolean with false value.
Definition: platform.h:308
cutlass::platform::unique_ptr::get_deleter
Deleter const & get_deleter() const noexcept
Returns the deleter object.
Definition: platform.h:786
cutlass::platform::remove_cv
std::remove_cv
Definition: platform.h:400
cutlass::platform::integral_constant::operator()
CUTLASS_HOST_DEVICE const value_type operator()() const
Definition: platform.h:294
cutlass::platform::unique_ptr::~unique_ptr
~unique_ptr()
Definition: platform.h:755
cutlass::platform::greater::operator()
CUTLASS_HOST_DEVICE constexpr bool operator()(const T &lhs, const T &rhs) const
Definition: platform.h:198
cutlass::platform::__align__
struct __align__(1) aligned_chunk< 1 >
Definition: platform.h:659
cutlass::platform::remove_const< const T >::type
T type
Definition: platform.h:383
cutlass::platform::remove_volatile::type
T type
Definition: platform.h:389
cutlass::platform::is_integral
std::is_integral
Definition: platform.h:492
cutlass::platform::is_trivially_copyable
Definition: platform.h:567
cutlass::platform::is_base_of_helper::dummy
Definition: platform.h:433
cutlass::platform::integral_constant::type
integral_constant< value_t, V > type
Definition: platform.h:290
cutlass::platform::is_arithmetic
std::is_arithmetic
Definition: platform.h:531
cutlass::platform::alignment_of::pad::byte
char byte
Definition: platform.h:587
cutlass::platform::integral_constant
std::integral_constant
Definition: platform.h:282
cutlass::platform::is_base_of
std::is_base_of
Definition: platform.h:448
cutlass::platform::enable_if::type
T type
Definition: platform.h:342
nullptr
#define nullptr
nullptr
Definition: platform.h:144
cutlass::platform::is_volatile
std::is_volatile
Definition: platform.h:470
cutlass::platform::is_fundamental
std::is_fundamental
Definition: platform.h:536
cutlass::platform::plus
platform::plus
Definition: platform.h:183
cutlass::platform::enable_if
std::enable_if (true specialization)
Definition: platform.h:341
cutlass::platform::true_type
integral_constant< bool, true > true_type
The type used as a compile-time boolean with true value.
Definition: platform.h:305
cutlass::platform::default_delete::operator()
void operator()(T *ptr) const
Definition: platform.h:727
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::platform::unique_ptr::element_type
T element_type
Definition: platform.h:741
cutlass::platform::unique_ptr::get_deleter
Deleter & get_deleter() noexcept
Returns the deleter object.
Definition: platform.h:783
cutlass::platform::alignment_of
std::alignment_of
Definition: platform.h:584
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:209
cutlass::platform::remove_cv::type
remove_volatile< typename remove_const< T >::type >::type type
Definition: platform.h:401
cutlass::platform::conditional
std::conditional (true specialization)
Definition: platform.h:351
noexcept
#define noexcept
noexcept, constexpr
Definition: platform.h:134
cutlass::platform::unique_ptr::reset
void reset(pointer p=pointer()) noexcept
Replaces the managed object, deleting the old object.
Definition: platform.h:771
cutlass::platform::unique_ptr::operator*
T & operator*() const
Dereferences the unique_ptr.
Definition: platform.h:792
cutlass::platform::is_base_of_helper
Helper for std::is_base_of.
Definition: platform.h:428
cutlass::platform::remove_const
std::remove_const (non-const specialization)
Definition: platform.h:376
cutlass::platform::plus::operator()
CUTLASS_HOST_DEVICE constexpr T operator()(const T &lhs, const T &rhs) const
Definition: platform.h:184
cutlass::platform::less::operator()
CUTLASS_HOST_DEVICE constexpr bool operator()(const T &lhs, const T &rhs) const
Definition: platform.h:190
cutlass::platform::aligned_chunk
Definition: platform.h:657
cutlass::platform::is_base_of_helper::check
static CUTLASS_HOST_DEVICE yes check(DerivedT *, T)
cutlass::platform::remove_volatile< volatile T >::type
T type
Definition: platform.h:395
cutlass::platform::operator!=
CUTLASS_HOST_DEVICE bool operator!=(complex< T > const &lhs, complex< T > const &rhs)
Inequality operator.
Definition: complex.h:232
cutlass::platform::unique_ptr::swap
void swap(unique_ptr &other) noexcept
Swaps the managed objects with *this and another unique_ptr.
Definition: platform.h:780
cutlass::platform::is_base_of_helper::value
static const bool value
Definition: platform.h:443
cutlass::platform::operator>=
CUTLASS_HOST_DEVICE bool operator>=(Pair< T1, T2 > const &lhs, Pair< T1, T2 > const &rhs)
Lexical comparison.
Definition: pair.h:117
cutlass::platform::aligned_storage::type
aligned_chunk< Align > type[Len/sizeof(aligned_chunk< Align >)]
Definition: platform.h:714
cutlass::platform::aligned_storage
std::aligned_storage
Definition: platform.h:713
cutlass::platform::remove_volatile
std::remove_volatile (non-volatile specialization)
Definition: platform.h:388
cutlass::platform::unique_ptr::unique_ptr
unique_ptr(pointer p)
Definition: platform.h:753
cutlass::platform::is_base_of_helper::yes
char(& yes)[1]
Definition: platform.h:429
cutlass::platform::unique_ptr::release
pointer release() noexcept
Releases ownership of the managed object, if any.
Definition: platform.h:764
cutlass.h
Basic include for CUTLASS macros.
cutlass::platform::bool_constant
std::bool_constant
Definition: platform.h:314
cutlass::platform::is_base_of_helper::no
char(& no)[2]
Definition: platform.h:430