diff options
Diffstat (limited to 'cutlass-example/cutlass/util/platform.h')
| -rw-r--r-- | cutlass-example/cutlass/util/platform.h | 801 |
1 files changed, 0 insertions, 801 deletions
diff --git a/cutlass-example/cutlass/util/platform.h b/cutlass-example/cutlass/util/platform.h deleted file mode 100644 index 2a44c10..0000000 --- a/cutlass-example/cutlass/util/platform.h +++ /dev/null @@ -1,801 +0,0 @@ -/*************************************************************************************************** - * Copyright (c) 2017-2018, 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 - -/** - * \file - * \brief C++ features that may be otherwise unimplemented for CUDA device functions. - * - * This file has three components: - * - * (1) Macros: - * - Empty macro defines for C++ keywords not supported by the current - * version of C++. These simply allow compilation to proceed (but do - * not provide the added semantics). - * - \p noexcept - * - \p constexpr - * - \p nullptr - * - \p static_assert - * - * - Macro functions that we need in constant expressions because the - * C++ equivalents require constexpr compiler support. These are - * prefixed with \p __NV_STD_* - * - \p __NV_STD_MAX - * - \p __NV_STD_MIN - * - * (2) Re-implementations of STL functions and types: - * - C++ features that need the \p __device__ annotation. These are - * placed into the \p platform namespace. - * - \p plus - * - \p less - * - \p greater - * - \p min - * - \p max - * - \p methods on std::pair (==, !=, <, <=, >, >=, and make_pair()) - * - * (3) Stop-gap implementations of unsupported STL functions and types: - * - STL functions and types defined by C++ 11/14/17/etc. that are not - * provided by the current version of C++. These are placed into the - * \p platform namespace - * - \p integral_constant - * - \p nullptr_t - * - \p true_type - * - \p false_type - * - \p bool_constant - * - \p enable_if - * - \p conditional - * - \p is_same - * - \p is_base_of - * - \p remove_const - * - \p remove_volatile - * - \p remove_cv - * - \p is_volatile - * - \p is_pointer - * - \p is_void - * - \p is_integral - * - \p is_floating_point - * - \p is_arithmetic - * - \p is_fundamental - * - \p is_trivially_copyable - * - \p alignment_of - * - \p aligned_storage - * - * (4) Functions and types that are STL-like (but aren't in the STL): - * - \p TODO: min and max functors? - * - * The idea is that, as we drop support for older compilers, we can simply #define - * the \p __NV_STD_XYZ macros and \p platform namespace to alias their C++ - * counterparts (or trivially find-and-replace their occurrences in code text). - */ - -//----------------------------------------------------------------------------- -// 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 -/****************************************************************************** - * Macros - ******************************************************************************/ -//----------------------------------------------------------------------------- -// Keywords -//----------------------------------------------------------------------------- - -/// noexcept, constexpr -#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1900)) -#ifndef noexcept -#define noexcept -#endif -#ifndef constexpr -#define constexpr -#endif -#endif - -/// nullptr -#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1310)) -#ifndef nullptr -#define nullptr 0 -#endif -#endif - -/// static_assert -#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 -//----------------------------------------------------------------------------- - -/// Select maximum(a, b) -#ifndef __NV_STD_MAX -#define __NV_STD_MAX(a, b) (((b) > (a)) ? (b) : (a)) -#endif - -/// Select minimum(a, b) -#ifndef __NV_STD_MIN -#define __NV_STD_MIN(a, b) (((b) < (a)) ? (b) : (a)) -#endif - -/****************************************************************************** - * Re-implementations - ******************************************************************************/ -namespace cutlass { -namespace platform { - -//----------------------------------------------------------------------------- -// Arithmetic operations, comparisons <functional> -//----------------------------------------------------------------------------- - -/// platform::plus -template <typename T> -struct plus { - CUTLASS_HOST_DEVICE constexpr T operator()(const T& lhs, const T& rhs) const { return lhs + rhs; } -}; - -/// std::less -template <typename T> -struct less { - CUTLASS_HOST_DEVICE constexpr bool operator()(const T& lhs, const T& rhs) const { - return lhs < rhs; - } -}; - -/// std::greater -template <typename T> -struct greater { - CUTLASS_HOST_DEVICE constexpr bool operator()(const T& lhs, const T& rhs) const { - return lhs > rhs; - } -}; - -//----------------------------------------------------------------------------- -// Minimum/maximum operations <algorithm> -//----------------------------------------------------------------------------- - -/// std::min -template <typename T> -CUTLASS_HOST_DEVICE constexpr const T& min(const T& a, const T& b) { - return (b < a) ? b : a; -} - -/// std::max -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500)) - -/// std::integral_constant -template <typename value_t, value_t V> -struct integral_constant; - -/// std::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 - -/// The type used as a compile-time boolean with true value. -typedef integral_constant<bool, true> true_type; - -/// The type used as a compile-time boolean with false value. -typedef integral_constant<bool, false> false_type; - -#if (!defined(_MSC_VER) && (__cplusplus <= 201402L)) || (defined(_MSC_VER) && (_MSC_VER < 1900)) - -/// std::bool_constant -template <bool V> -struct bool_constant : platform::integral_constant<bool, V> {}; - -#else - -using std::bool_constant; - -#endif - -#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1700)) - -/// std::nullptr_t -struct nullptr_t {}; - -#else - -using std::nullptr_t; - -#endif - -//----------------------------------------------------------------------------- -// Conditional metaprogramming <type_traits> -//----------------------------------------------------------------------------- - -#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600)) - -/// std::enable_if (true specialization) -template <bool C, typename T = void> -struct enable_if { - typedef T type; -}; - -/// std::enable_if (false specialization) -template <typename T> -struct enable_if<false, T> {}; - -/// std::conditional (true specialization) -template <bool B, class T, class F> -struct conditional { - typedef T type; -}; - -/// std::conditional (false specialization) -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500)) - -/// std::remove_const (non-const specialization) -template <typename T> -struct remove_const { - typedef T type; -}; - -/// std::remove_const (const specialization) -template <typename T> -struct remove_const<const T> { - typedef T type; -}; - -/// std::remove_volatile (non-volatile specialization) -template <typename T> -struct remove_volatile { - typedef T type; -}; - -/// std::remove_volatile (volatile specialization) -template <typename T> -struct remove_volatile<volatile T> { - typedef T type; -}; - -/// std::remove_cv -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500)) - -/// std::is_same (false specialization) -template <typename A, typename B> -struct is_same : false_type {}; - -/// std::is_same (true specialization) -template <typename A> -struct is_same<A, A> : true_type {}; - -/// Helper for std::is_base_of -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); -}; - -/// std::is_base_of -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500)) - -/// std::is_volatile -template <typename T> -struct is_volatile : false_type {}; -template <typename T> -struct is_volatile<volatile T> : true_type {}; - -/// Helper for std::is_pointer (false specialization) -template <typename T> -struct is_pointer_helper : false_type {}; - -/// Helper for std::is_pointer (true specialization) -template <typename T> -struct is_pointer_helper<T*> : true_type {}; - -/// std::is_pointer -template <typename T> -struct is_pointer : is_pointer_helper<typename remove_cv<T>::type> {}; - -/// std::is_void -template <typename T> -struct is_void : is_same<void, typename remove_cv<T>::type> {}; - -/// std::is_integral -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> {}; - -/// std::is_floating_point -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)> {}; - -/// std::is_arithmetic -template <typename T> -struct is_arithmetic - : integral_constant<bool, (is_integral<T>::value || is_floating_point<T>::value)> {}; - -/// std::is_fundamental -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800)) || \ - (defined(__GNUG__) && (__GNUC__ < 5)) - -/** - * std::is_trivially_copyable - * - * This implementation only evaluates true if T is fundamental or pointer - * - * Without help from partial template specializations provided by the user for - * a specific class or struct, this trait will never report that the specified - * class or struct is trivially-copyable ; this is always safe, - * if possibly sub-optimal. - */ -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(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500)) - -/// std::alignment_of -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(_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]; -}; - -/// std::aligned_storage -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__) -/// Default deleter -template <typename T> -struct default_delete { - void operator()(T* ptr) const { delete ptr; } -}; - -/// Partial specialization for deleting array types -template <typename T> -struct default_delete<T[]> { - void operator()(T* ptr) const { delete[] ptr; } -}; - -/// std::unique_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 to memory - pointer _ptr; - - /// Deleter - deleter_type _deleter; - - public: - unique_ptr() : _ptr(nullptr) {} - unique_ptr(pointer p) : _ptr(p) {} - - ~unique_ptr() { - if (_ptr) { - _deleter(_ptr); - } - } - /// Returns a pointer to the managed object or nullptr if no object is owned. - pointer get() const noexcept { return _ptr; } - - /// Releases ownership of the managed object, if any - pointer release() noexcept { - pointer p(_ptr); - _ptr = nullptr; - return p; - } - - /// Replaces the managed object, deleting the old object. - void reset(pointer p = pointer()) noexcept { - pointer old_ptr = _ptr; - _ptr = p; - if (old_ptr != nullptr) { - get_deleter()(old_ptr); - } - } - - /// Swaps the managed objects with *this and another unique_ptr - void swap(unique_ptr& other) noexcept { std::swap(_ptr, other._ptr); } - - /// Returns the deleter object - Deleter& get_deleter() noexcept { return _deleter; } - - /// Returns the deleter object - Deleter const& get_deleter() const noexcept { return _deleter; } - - /// Checks whether an object is owned - operator bool() const noexcept { return _ptr != nullptr; } - - /// Dereferences the unique_ptr - T& operator*() const { return *_ptr; } - - /// Returns a pointer to the managed object - pointer operator->() const noexcept { return _ptr; } - - /// Array access to managed object - T& operator[](size_t i) const { return _ptr[i]; } -}; - -/// Specializes the swap algorithm -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 |