1 files changed, 160 insertions, 0 deletions
diff --git a/cutlass-example/type_traits.h b/cutlass-example/type_traits.h
new file mode 100644
index 0000000..eabd67e
--- /dev/null
+++ b/cutlass-example/type_traits.h
@@ -0,0 +1,160 @@
+/***************************************************************************************************
+ * 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.
+ *
+ **************************************************************************************************/
+/*! \file
+    \brief Type traits for common CUDA types
+*/
+
+#pragma once
+
+#include <cuda_fp16.h>
+#include <stdint.h>
+
+#include "half.h"
+
+namespace cutlass {
+struct half_t;
+
+template <typename T>
+struct TypeTraits;
+
+template <>
+struct TypeTraits<int8_t> {
+//  static cudaDataType_t const cublas_type = CUDA_R_8I;
+  typedef int8_t host_type;
+  typedef int8_t device_type;
+  typedef int8_t integer_type;
+  typedef uint8_t unsigned_type;
+  static inline int8_t remove_negative_zero(int8_t x) { return x; }
+  static inline int to_print(int8_t x) { return (int)x; }
+};
+
+template <>
+struct TypeTraits<uint8_t> {
+//  static cudaDataType_t const cublas_type = CUDA_R_8I;
+  typedef uint8_t host_type;
+  typedef uint8_t device_type;
+  typedef uint8_t integer_type;
+  typedef uint8_t unsigned_type;
+  static inline uint8_t remove_negative_zero(uint8_t x) { return x; }
+  static inline uint32_t to_print(uint8_t x) { return (uint32_t)x; }
+};
+
+template <>
+struct TypeTraits<int> {
+//  static cudaDataType_t const cublas_type = CUDA_R_32I;
+  typedef int host_type;
+  typedef int device_type;
+  typedef int32_t integer_type;
+  typedef uint32_t unsigned_type;
+  static inline int32_t remove_negative_zero(int32_t x) { return x; }
+  static inline int to_print(int x) { return x; }
+};
+
+template <>
+struct TypeTraits<unsigned> {
+//  static cudaDataType_t const cublas_type = CUDA_R_32I;
+  typedef unsigned host_type;
+  typedef unsigned device_type;
+  typedef uint32_t integer_type;
+  typedef uint32_t unsigned_type;
+  static inline uint32_t remove_negative_zero(uint32_t x) { return x; }
+  static inline uint32_t to_print(uint32_t x) { return x; }
+};
+
+template <>
+struct TypeTraits<half> {
+//  static cudaDataType_t const cublas_type = CUDA_R_16F;
+  typedef half_t host_type;
+  typedef half device_type;
+  typedef int16_t integer_type;
+  typedef uint16_t unsigned_type;
+  static inline half remove_negative_zero(half x) {
+    integer_type h_int = reinterpret_cast<integer_type const&>(x);
+    if (h_int == 0x8000) {
+      h_int = 0;
+    }
+    x = reinterpret_cast<half const&>(h_int);
+    return x;
+  }
+  static inline half to_print(half x) { return x; }
+};
+
+template <>
+struct TypeTraits<int64_t> {
+//  static cudaDataType_t const cublas_type = CUDA_R_8I;
+  typedef int64_t host_type;
+  typedef int64_t device_type;
+  typedef int64_t integer_type;
+  typedef uint64_t unsigned_type;
+  static inline int64_t remove_negative_zero(int64_t x) { return x; }
+  static inline int64_t to_print(int64_t x) { return x; }
+};
+
+template <>
+struct TypeTraits<uint64_t> {
+//  static cudaDataType_t const cublas_type = CUDA_R_8I;
+  typedef uint64_t host_type;
+  typedef uint64_t device_type;
+  typedef uint64_t integer_type;
+  typedef uint64_t unsigned_type;
+  static inline uint64_t remove_negative_zero(uint64_t x) { return x; }
+  static inline uint64_t to_print(uint64_t x) { return x; }
+};
+
+template <>
+struct TypeTraits<cutlass::half_t> {
+//  static cudaDataType_t const cublas_type = CUDA_R_16F;
+  typedef half_t host_type;
+  typedef half device_type;
+  typedef int16_t integer_type;
+  typedef uint16_t unsigned_type;
+  static inline half_t remove_negative_zero(half_t x) {
+    return (x.raw() == 0x8000 ? half_t::bitcast(0) : x);
+  }
+  static inline half_t to_print(half_t x) { return x; }
+};
+
+template <>
+struct TypeTraits<float> {
+//  static cudaDataType_t const cublas_type = CUDA_R_32F;
+  typedef float host_type;
+  typedef float device_type;
+  typedef int32_t integer_type;
+  typedef uint32_t unsigned_type;
+  static inline float remove_negative_zero(float x) { return x == -0.f ? 0.f : x; }
+  static inline float to_print(float x) { return x; }
+};
+
+template <>
+struct TypeTraits<double> {
+//  static cudaDataType_t const cublas_type = CUDA_R_64F;
+  typedef double host_type;
+  typedef double device_type;
+  typedef int64_t integer_type;
+  typedef uint64_t unsigned_type;
+  static inline double remove_negative_zero(double x) { return x == -0.0 ? 0.0 : x; }
+  static inline double to_print(double x) { return x; }
+};
+}  // namespace cutlass