diff options
Diffstat (limited to 'cutlass-example/cutlass/util')
| -rw-r--r-- | cutlass-example/cutlass/util/cutlass_math.h | 131 | ||||
| -rw-r--r-- | cutlass-example/cutlass/util/debug.h | 122 | ||||
| -rw-r--r-- | cutlass-example/cutlass/util/platform.h | 801 |
3 files changed, 1054 insertions, 0 deletions
diff --git a/cutlass-example/cutlass/util/cutlass_math.h b/cutlass-example/cutlass/util/cutlass_math.h new file mode 100644 index 0000000..0ecdc43 --- /dev/null +++ b/cutlass-example/cutlass/util/cutlass_math.h @@ -0,0 +1,131 @@ +/*************************************************************************************************** + * 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 Math utilities + */ + +#include <cutlass/util/platform.h> + +namespace cutlass { + +/****************************************************************************** + * Static math utilities + ******************************************************************************/ + +/** + * Statically determine if N is a power-of-two + */ +template <int N> +struct is_pow2 : platform::integral_constant<bool, (N & (N - 1)) == 0> {}; + +/** + * Statically determine log2(N), rounded down + */ +template <int N, int CurrentVal = N, int Count = 0> +struct log2_down { + /// Static logarithm value + enum { value = log2_down<N, (CurrentVal >> 1), Count + 1>::value }; +}; + +// Base case +template <int N, int Count> +struct log2_down<N, 1, Count> { + enum { value = Count }; +}; + +/** + * Statically determine log2(N), rounded up + */ +template <int N, int CurrentVal = N, int Count = 0> +struct log2_up { + /// Static logarithm value + enum { value = log2_up<N, (CurrentVal >> 1), Count + 1>::value }; +}; + +// Base case +template <int N, int Count> +struct log2_up<N, 1, Count> { + enum { value = ((1 << Count) < N) ? Count + 1 : Count }; +}; + +/** + * Statically estimate sqrt(N) to the nearest power-of-two + */ +template <int N> +struct sqrt_est { + enum { value = 1 << (log2_up<N>::value / 2) }; +}; + +/** + * For performing a constant-division with a compile-time assertion that the + * Divisor evenly-divides the Dividend. + */ +template <int Dividend, int Divisor> +struct divide_assert { + enum { value = Dividend / Divisor }; + + static_assert((Dividend % Divisor == 0), "Not an even multiple"); +}; + +/****************************************************************************** + * Rounding + ******************************************************************************/ + +/** + * Round dividend up to the nearest multiple of divisor + */ +template <typename dividend_t, typename divisor_t> +CUTLASS_HOST_DEVICE dividend_t round_nearest(dividend_t dividend, divisor_t divisor) { + return ((dividend + divisor - 1) / divisor) * divisor; +} + +/** + * Greatest common divisor + */ +template <typename value_t> +CUTLASS_HOST_DEVICE value_t gcd(value_t a, value_t b) { + for (;;) { + if (a == 0) return b; + b %= a; + if (b == 0) return a; + a %= b; + } +} + +/** + * Least common multiple + */ +template <typename value_t> +CUTLASS_HOST_DEVICE value_t lcm(value_t a, value_t b) { + value_t temp = gcd(a, b); + + return temp ? (a / temp * b) : 0; +} + +} // namespace cutlass diff --git a/cutlass-example/cutlass/util/debug.h b/cutlass-example/cutlass/util/debug.h new file mode 100644 index 0000000..6055e3f --- /dev/null +++ b/cutlass-example/cutlass/util/debug.h @@ -0,0 +1,122 @@ +/*************************************************************************************************** + * 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 Debugging and logging functionality + */ + +#include <stdio.h> + +namespace cutlass { + +/****************************************************************************** + * Debug and logging macros + ******************************************************************************/ + +/** + * Formats and prints the given message to stdout + */ +#if !defined(CUDA_LOG) +#if !defined(__CUDA_ARCH__) +#define CUDA_LOG(format, ...) printf(format, __VA_ARGS__) +#else +#define CUDA_LOG(format, ...) \ + printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format, \ + blockIdx.x, \ + blockIdx.y, \ + blockIdx.z, \ + threadIdx.x, \ + threadIdx.y, \ + threadIdx.z, \ + __VA_ARGS__); +#endif +#endif + +/** + * Formats and prints the given message to stdout only if DEBUG is defined + */ +#if !defined(CUDA_LOG_DEBUG) +#ifdef DEBUG +#define CUDA_LOG_DEBUG(format, ...) CUDA_LOG(format, __VA_ARGS__) +#else +#define CUDA_LOG_DEBUG(format, ...) +#endif +#endif + +/** + * \brief The corresponding error message is printed to \p stderr (or \p stdout in device code) + * along with the supplied source context. + * + * \return The CUDA error. + */ +__host__ CUTLASS_DEVICE cudaError_t cuda_perror_impl(cudaError_t error, + const char* filename, + int line) { + (void)filename; + (void)line; + if (error) { +#if !defined(__CUDA_ARCH__) + fprintf( + stderr, "CUDA error %d [%s, %d]: %s\n", error, filename, line, cudaGetErrorString(error)); + fflush(stderr); +#else + printf("CUDA error %d [%s, %d]\n", error, filename, line); +#endif + } + return error; +} + +/** + * \brief Perror macro + */ +#ifndef CUDA_PERROR +#define CUDA_PERROR(e) cuda_perror_impl((cudaError_t)(e), __FILE__, __LINE__) +#endif + +/** + * \brief Perror macro with exit + */ +#ifndef CUDA_PERROR_EXIT +#define CUDA_PERROR_EXIT(e) \ + if (cuda_perror_impl((cudaError_t)(e), __FILE__, __LINE__)) { \ + exit(1); \ + } +#endif + +/** + * \brief Perror macro only if DEBUG is defined + */ +#ifndef CUDA_PERROR_DEBUG +#ifdef DEBUG +#define CUDA_PERROR_DEBUG(e) CUDA_PERROR(e) +#else +#define CUDA_PERROR_DEBUG(e) (e) +#endif +#endif + +} // namespace cutlass diff --git a/cutlass-example/cutlass/util/platform.h b/cutlass-example/cutlass/util/platform.h new file mode 100644 index 0000000..2a44c10 --- /dev/null +++ b/cutlass-example/cutlass/util/platform.h @@ -0,0 +1,801 @@ +/*************************************************************************************************** + * 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 |