platform.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification, are permitted
* provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright notice, this list of
* conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright notice, this list of
* conditions and the following disclaimer in the documentation and/or other materials
* provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
* to endorse or promote products derived from this software without specific prior written
* permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
* IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
* FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
* STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
#pragma once
//-----------------------------------------------------------------------------
// Dependencies
//-----------------------------------------------------------------------------
#include <stdint.h>
#if !defined(__CUDACC_RTC__)
//-----------------------------------------------------------------------------
// Include STL files that platform provides functionality for
//-----------------------------------------------------------------------------
#include <algorithm> // Minimum/maximum operations
#include <cstddef> // nullptr_t
#include <functional> // Arithmetic operations
#include <utility> // For methods on std::pair
#if (!defined(_MSC_VER) && (__cplusplus >= 201103L)) || (defined(_MSC_VER) && (_MS_VER >= 1500))
#include <type_traits> // For integral constants, conditional metaprogramming, and type traits
#endif
#include "cutlass/cutlass.h"
#endif
//-----------------------------------------------------------------------------
// OS
//-----------------------------------------------------------------------------
#if defined(WIN32) || defined(_WIN32) || defined(__WIN32) && !defined(__CYGWIN__)
#define CUTLASS_OS_WINDOWS
#endif
/******************************************************************************
* Macros
******************************************************************************/
//-----------------------------------------------------------------------------
// Keywords
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
#ifndef noexcept
#define noexcept
#endif
#ifndef constexpr
#define constexpr
#endif
#endif
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1310))
#ifndef nullptr
#define nullptr 0
#endif
#endif
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600))
#ifndef static_assert
#define __platform_cat_(a, b) a##b
#define __platform_cat(a, b) __platform_cat_(a, b)
#define static_assert(__e, __m) typedef int __platform_cat(AsSeRt, __LINE__)[(__e) ? 1 : -1]
#endif
#endif
//-----------------------------------------------------------------------------
// Functions
//-----------------------------------------------------------------------------
#ifndef __NV_STD_MAX
#define __NV_STD_MAX(a, b) (((b) > (a)) ? (b) : (a))
#endif
#ifndef __NV_STD_MIN
#define __NV_STD_MIN(a, b) (((b) < (a)) ? (b) : (a))
#endif
/******************************************************************************
* Re-implementations
******************************************************************************/
namespace cutlass {
namespace platform {
//-----------------------------------------------------------------------------
// Minimum/maximum operations <algorithm>
//-----------------------------------------------------------------------------
template <typename T>
CUTLASS_HOST_DEVICE constexpr const T& min(const T& a, const T& b) {
return (b < a) ? b : a;
}
template <typename T>
CUTLASS_HOST_DEVICE constexpr const T& max(const T& a, const T& b) {
return (a < b) ? b : a;
}
#if !defined(__CUDACC_RTC__)
//-----------------------------------------------------------------------------
// Methods on std::pair
//-----------------------------------------------------------------------------
using std::pair;
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator==(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return (lhs.first == rhs.first) && (lhs.second == rhs.second);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator!=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return (lhs.first != rhs.first) && (lhs.second != rhs.second);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator<(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return (lhs.first < rhs.first) ? true : (rhs.first < lhs.first) ? false
: (lhs.second < rhs.second);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator<=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return !(rhs < lhs);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator>(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return (rhs < lhs);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE constexpr bool operator>=(const pair<T1, T2>& lhs, const pair<T1, T2>& rhs) {
return !(lhs < rhs);
}
template <class T1, class T2>
CUTLASS_HOST_DEVICE std::pair<T1, T2> make_pair(T1 t, T2 u) {
std::pair<T1, T2> retval;
retval.first = t;
retval.second = u;
return retval;
}
#endif
} // namespace platform
/******************************************************************************
* Implementations of C++ 11/14/17/... STL features
******************************************************************************/
namespace platform {
//-----------------------------------------------------------------------------
// Integral constant helper types <type_traits>
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
template <typename value_t, value_t V>
struct integral_constant;
template <typename value_t, value_t V>
struct integral_constant {
static const value_t value = V;
typedef value_t value_type;
typedef integral_constant<value_t, V> type;
CUTLASS_HOST_DEVICE operator value_type() const { return value; }
CUTLASS_HOST_DEVICE const value_type operator()() const { return value; }
};
#else
using std::integral_constant;
using std::pair;
#endif
typedef integral_constant<bool, true> true_type;
typedef integral_constant<bool, false> false_type;
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus <= 201402L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
template <bool V>
struct bool_constant : platform::integral_constant<bool, V> {};
#else
using std::bool_constant;
#endif
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1700))
struct nullptr_t {};
#else
using std::nullptr_t;
#endif
//-----------------------------------------------------------------------------
// Conditional metaprogramming <type_traits>
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600))
template <bool C, typename T = void>
struct enable_if {
typedef T type;
};
template <typename T>
struct enable_if<false, T> {};
template <bool B, class T, class F>
struct conditional {
typedef T type;
};
template <class T, class F>
struct conditional<false, T, F> {
typedef F type;
};
#else
using std::enable_if;
using std::conditional;
#endif
//-----------------------------------------------------------------------------
// Const/volatility specifiers <type_traits>
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
template <typename T>
struct remove_const {
typedef T type;
};
template <typename T>
struct remove_const<const T> {
typedef T type;
};
template <typename T>
struct remove_volatile {
typedef T type;
};
template <typename T>
struct remove_volatile<volatile T> {
typedef T type;
};
template <typename T>
struct remove_cv {
typedef typename remove_volatile<typename remove_const<T>::type>::type type;
};
#else
using std::remove_const;
using std::remove_volatile;
using std::remove_cv;
#endif
//-----------------------------------------------------------------------------
// Type relationships <type_traits>
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
template <typename A, typename B>
struct is_same : false_type {};
template <typename A>
struct is_same<A, A> : true_type {};
template <typename BaseT, typename DerivedT>
struct is_base_of_helper {
typedef char (&yes)[1];
typedef char (&no)[2];
template <typename B, typename D>
struct dummy {
CUTLASS_HOST_DEVICE operator B*() const;
CUTLASS_HOST_DEVICE operator D*();
};
template <typename T>
CUTLASS_HOST_DEVICE static yes check(DerivedT*, T);
CUTLASS_HOST_DEVICE static no check(BaseT*, int);
static const bool value = sizeof(check(dummy<BaseT, DerivedT>(), int())) == sizeof(yes);
};
template <typename BaseT, typename DerivedT>
struct is_base_of
: integral_constant<bool,
(is_base_of_helper<typename remove_cv<BaseT>::type,
typename remove_cv<DerivedT>::type>::value) ||
(is_same<typename remove_cv<BaseT>::type,
typename remove_cv<DerivedT>::type>::value)> {};
#else
using std::is_same;
using std::is_base_of;
#endif
//-----------------------------------------------------------------------------
// Type properties <type_traits>
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
template <typename T>
struct is_volatile : false_type {};
template <typename T>
struct is_volatile<volatile T> : true_type {};
template <typename T>
struct is_pointer_helper : false_type {};
template <typename T>
struct is_pointer_helper<T*> : true_type {};
template <typename T>
struct is_pointer : is_pointer_helper<typename remove_cv<T>::type> {};
template <typename T>
struct is_void : is_same<void, typename remove_cv<T>::type> {};
template <typename T>
struct is_integral : false_type {};
template <>
struct is_integral<char> : true_type {};
template <>
struct is_integral<signed char> : true_type {};
template <>
struct is_integral<unsigned char> : true_type {};
template <>
struct is_integral<short> : true_type {};
template <>
struct is_integral<unsigned short> : true_type {};
template <>
struct is_integral<int> : true_type {};
template <>
struct is_integral<unsigned int> : true_type {};
template <>
struct is_integral<long> : true_type {};
template <>
struct is_integral<unsigned long> : true_type {};
template <>
struct is_integral<long long> : true_type {};
template <>
struct is_integral<unsigned long long> : true_type {};
template <typename T>
struct is_integral<volatile T> : is_integral<T> {};
template <typename T>
struct is_integral<const T> : is_integral<T> {};
template <typename T>
struct is_integral<const volatile T> : is_integral<T> {};
template <typename T>
struct is_floating_point
: integral_constant<bool,
(is_same<float, typename remove_cv<T>::type>::value ||
is_same<double, typename remove_cv<T>::type>::value)> {};
template <typename T>
struct is_arithmetic
: integral_constant<bool, (is_integral<T>::value || is_floating_point<T>::value)> {};
template <typename T>
struct is_fundamental
: integral_constant<bool,
(is_arithmetic<T>::value || is_void<T>::value ||
is_same<nullptr_t, typename remove_cv<T>::type>::value)> {};
#else
using std::is_volatile;
using std::is_pointer;
using std::is_void;
using std::is_integral;
using std::is_floating_point;
using std::is_arithmetic;
using std::is_fundamental;
#endif
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800)) || \
(defined(__GNUG__) && (__GNUC__ < 5))
template <typename T>
struct is_trivially_copyable
: integral_constant<bool, (is_fundamental<T>::value || is_pointer<T>::value)> {};
#else
using std::is_trivially_copyable;
#endif
//-----------------------------------------------------------------------------
// Alignment and layout utilities
//-----------------------------------------------------------------------------
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
template <typename value_t>
struct alignment_of {
struct pad {
value_t val;
char byte;
};
enum { value = sizeof(pad) - sizeof(value_t) };
};
#else
template <typename value_t>
struct alignment_of : std::alignment_of<value_t> {};
#endif
/* 16B specializations where 32-bit Win32 host compiler disagrees with device compiler */
template <>
struct alignment_of<int4> {
enum { value = 16 };
};
template <>
struct alignment_of<uint4> {
enum { value = 16 };
};
template <>
struct alignment_of<float4> {
enum { value = 16 };
};
template <>
struct alignment_of<long4> {
enum { value = 16 };
};
template <>
struct alignment_of<ulong4> {
enum { value = 16 };
};
template <>
struct alignment_of<longlong2> {
enum { value = 16 };
};
template <>
struct alignment_of<ulonglong2> {
enum { value = 16 };
};
template <>
struct alignment_of<double2> {
enum { value = 16 };
};
template <>
struct alignment_of<longlong4> {
enum { value = 16 };
};
template <>
struct alignment_of<ulonglong4> {
enum { value = 16 };
};
template <>
struct alignment_of<double4> {
enum { value = 16 };
};
// Specializations for volatile/const qualified types
template <typename value_t>
struct alignment_of<volatile value_t> : alignment_of<value_t> {};
template <typename value_t>
struct alignment_of<const value_t> : alignment_of<value_t> {};
template <typename value_t>
struct alignment_of<const volatile value_t> : alignment_of<value_t> {};
#if defined(__CUDACC_RTC__) || (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800))
template <size_t Align>
struct aligned_chunk;
template <>
struct __align__(1) aligned_chunk<1> {
uint8_t buff;
};
template <>
struct __align__(2) aligned_chunk<2> {
uint16_t buff;
};
template <>
struct __align__(4) aligned_chunk<4> {
uint32_t buff;
};
template <>
struct __align__(8) aligned_chunk<8> {
uint32_t buff[2];
};
template <>
struct __align__(16) aligned_chunk<16> {
uint32_t buff[4];
};
template <>
struct __align__(32) aligned_chunk<32> {
uint32_t buff[8];
};
template <>
struct __align__(64) aligned_chunk<64> {
uint32_t buff[16];
};
template <>
struct __align__(128) aligned_chunk<128> {
uint32_t buff[32];
};
template <>
struct __align__(256) aligned_chunk<256> {
uint32_t buff[64];
};
template <>
struct __align__(512) aligned_chunk<512> {
uint32_t buff[128];
};
template <>
struct __align__(1024) aligned_chunk<1024> {
uint32_t buff[256];
};
template <>
struct __align__(2048) aligned_chunk<2048> {
uint32_t buff[512];
};
template <>
struct __align__(4096) aligned_chunk<4096> {
uint32_t buff[1024];
};
template <size_t Len, size_t Align>
struct aligned_storage {
typedef aligned_chunk<Align> type[Len / sizeof(aligned_chunk<Align>)];
};
#else
using std::aligned_storage;
#endif
#if !defined(__CUDACC_RTC__)
template <typename T>
struct default_delete {
void operator()(T* ptr) const { delete ptr; }
};
template <typename T>
struct default_delete<T[]> {
void operator()(T* ptr) const { delete[] ptr; }
};
template <class T, class Deleter = default_delete<T> >
class unique_ptr {
public:
typedef T* pointer;
typedef T element_type;
typedef Deleter deleter_type;
private:
pointer _ptr;
deleter_type _deleter;
public:
unique_ptr() : _ptr(nullptr) {}
unique_ptr(pointer p) : _ptr(p) {}
~unique_ptr() {
if (_ptr) {
_deleter(_ptr);
}
}
pointer get() const noexcept { return _ptr; }
pointer release() noexcept {
pointer p(_ptr);
_ptr = nullptr;
return p;
}
void reset(pointer p = pointer()) noexcept {
pointer old_ptr = _ptr;
_ptr = p;
if (old_ptr != nullptr) {
get_deleter()(old_ptr);
}
}
void swap(unique_ptr& other) noexcept { std::swap(_ptr, other._ptr); }
Deleter& get_deleter() noexcept { return _deleter; }
Deleter const& get_deleter() const noexcept { return _deleter; }
operator bool() const noexcept { return _ptr != nullptr; }
T& operator*() const { return *_ptr; }
pointer operator->() const noexcept { return _ptr; }
T& operator[](size_t i) const { return _ptr[i]; }
};
template <typename T, typename Deleter>
void swap(unique_ptr<T, Deleter>& lhs, unique_ptr<T, Deleter>& rhs) noexcept {
lhs.swap(rhs);
}
#endif
} // namespace platform
} // namespace cutlass
cutlass::platform::unique_ptr::operator[]
T & operator[](size_t i) const
Array access to managed object.
Definition: platform.h:772
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:189
cutlass
Definition: aligned_buffer.h:35
constexpr
#define constexpr
Definition: platform.h:137
cutlass::platform::nullptr_t
std::nullptr_t
Definition: platform.h:299
cutlass::platform::swap
void swap(unique_ptr< T, Deleter > &lhs, unique_ptr< T, Deleter > &rhs) noexcept
Specializes the swap algorithm.
Definition: platform.h:777
cutlass::platform::is_pointer_helper
Helper for std::is_pointer (false specialization)
Definition: platform.h:450
cutlass::platform::unique_ptr::deleter_type
Deleter deleter_type
Definition: platform.h:716
cutlass::platform::remove_const::type
T type
Definition: platform.h:351
cutlass::platform::alignment_of::pad::val
value_t val
Definition: platform.h:560
cutlass::platform::conditional::type
T type
Definition: platform.h:326
cutlass::platform::unique_ptr::pointer
T * pointer
Definition: platform.h:714
cutlass::platform::is_same
std::is_same (false specialization)
Definition: platform.h:394
cutlass::platform::is_pointer
std::is_pointer
Definition: platform.h:458
cutlass::platform::integral_constant::value_type
value_t value_type
Definition: platform.h:263
cutlass::platform::make_pair
CUTLASS_HOST_DEVICE std::pair< T1, T2 > make_pair(T1 t, T2 u)
Definition: platform.h:232
cutlass::platform::unique_ptr::unique_ptr
unique_ptr()
Definition: platform.h:726
cutlass::platform::default_delete< T[]>::operator()
void operator()(T *ptr) const
Definition: platform.h:707
cutlass::platform::operator==
CUTLASS_HOST_DEVICE constexpr bool operator==(const pair< T1, T2 > &lhs, const pair< T1, T2 > &rhs)
Definition: platform.h:201
cutlass::platform::is_void
std::is_void
Definition: platform.h:462
cutlass::platform::operator>=
CUTLASS_HOST_DEVICE constexpr bool operator>=(const pair< T1, T2 > &lhs, const pair< T1, T2 > &rhs)
Definition: platform.h:227
cutlass::platform::unique_ptr::operator->
pointer operator->() const noexcept
Returns a pointer to the managed object.
Definition: platform.h:769
cutlass::platform::operator>
CUTLASS_HOST_DEVICE constexpr bool operator>(const pair< T1, T2 > &lhs, const pair< T1, T2 > &rhs)
Definition: platform.h:222
cutlass::platform::default_delete
Default deleter.
Definition: platform.h:700
cutlass::platform::operator!=
CUTLASS_HOST_DEVICE constexpr bool operator!=(const pair< T1, T2 > &lhs, const pair< T1, T2 > &rhs)
Definition: platform.h:206
cutlass::platform::unique_ptr
std::unique_ptr
Definition: platform.h:712
cutlass::platform::alignment_of::pad
Definition: platform.h:559
cutlass::platform::is_floating_point
std::is_floating_point
Definition: platform.h:498
cutlass::platform::conditional< false, T, F >::type
F type
Definition: platform.h:332
cutlass::platform::false_type
integral_constant< bool, false > false_type
The type used as a compile-time boolean with false value.
Definition: platform.h:282
cutlass::platform::unique_ptr::get_deleter
Deleter const & get_deleter() const noexcept
Returns the deleter object.
Definition: platform.h:760
cutlass::platform::remove_cv
std::remove_cv
Definition: platform.h:374
cutlass::platform::unique_ptr::~unique_ptr
~unique_ptr()
Definition: platform.h:729
cutlass::platform::__align__
struct __align__(1) aligned_chunk< 1 >
Definition: platform.h:633
cutlass::platform::remove_const< const T >::type
T type
Definition: platform.h:357
cutlass::platform::remove_volatile::type
T type
Definition: platform.h:363
cutlass::platform::is_integral
std::is_integral
Definition: platform.h:466
cutlass::platform::is_trivially_copyable
Definition: platform.h:541
cutlass::platform::is_base_of_helper::dummy
Definition: platform.h:407
cutlass::platform::integral_constant::type
integral_constant< value_t, V > type
Definition: platform.h:264
cutlass::platform::is_arithmetic
std::is_arithmetic
Definition: platform.h:505
cutlass::platform::alignment_of::pad::byte
char byte
Definition: platform.h:561
cutlass::platform::integral_constant
std::integral_constant
Definition: platform.h:256
cutlass::platform::is_base_of
std::is_base_of
Definition: platform.h:422
cutlass::platform::enable_if::type
T type
Definition: platform.h:316
nullptr
#define nullptr
nullptr
Definition: platform.h:144
cutlass::platform::is_volatile
std::is_volatile
Definition: platform.h:444
cutlass::platform::is_fundamental
std::is_fundamental
Definition: platform.h:510
cutlass::platform::unique_ptr::operator*
T & operator*() const
Dereferences the unique_ptr.
Definition: platform.h:766
cutlass::platform::enable_if
std::enable_if (true specialization)
Definition: platform.h:315
cutlass::platform::true_type
integral_constant< bool, true > true_type
The type used as a compile-time boolean with true value.
Definition: platform.h:279
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::platform::unique_ptr::element_type
T element_type
Definition: platform.h:715
cutlass::platform::unique_ptr::get_deleter
Deleter & get_deleter() noexcept
Returns the deleter object.
Definition: platform.h:757
cutlass::platform::alignment_of
std::alignment_of
Definition: platform.h:558
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:183
cutlass::platform::remove_cv::type
remove_volatile< typename remove_const< T >::type >::type type
Definition: platform.h:375
cutlass::platform::conditional
std::conditional (true specialization)
Definition: platform.h:325
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:745
cutlass::platform::integral_constant::operator()
CUTLASS_HOST_DEVICE const value_type operator()() const
Definition: platform.h:268
cutlass::platform::is_base_of_helper
Helper for std::is_base_of.
Definition: platform.h:402
cutlass::platform::remove_const
std::remove_const (non-const specialization)
Definition: platform.h:350
cutlass::platform::aligned_chunk
Definition: platform.h:631
cutlass::platform::remove_volatile< volatile T >::type
T type
Definition: platform.h:369
cutlass::platform::unique_ptr::swap
void swap(unique_ptr &other) noexcept
Swaps the managed objects with *this and another unique_ptr.
Definition: platform.h:754
cutlass::platform::default_delete::operator()
void operator()(T *ptr) const
Definition: platform.h:701
cutlass::platform::aligned_storage
std::aligned_storage
Definition: platform.h:687
cutlass::platform::remove_volatile
std::remove_volatile (non-volatile specialization)
Definition: platform.h:362
cutlass::platform::unique_ptr::unique_ptr
unique_ptr(pointer p)
Definition: platform.h:727
cutlass::platform::unique_ptr::release
pointer release() noexcept
Releases ownership of the managed object, if any.
Definition: platform.h:738
cutlass.h
Basic include for CUTLASS.
cutlass::platform::bool_constant
std::bool_constant
Definition: platform.h:288