adding the cutlass examples

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