| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

type_traits.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without modification, are permitted
5 * provided that the following conditions are met:
6 * * Redistributions of source code must retain the above copyright notice, this list of
7 * conditions and the following disclaimer.
8 * * Redistributions in binary form must reproduce the above copyright notice, this list of
9 * conditions and the following disclaimer in the documentation and/or other materials
10 * provided with the distribution.
11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12 * to endorse or promote products derived from this software without specific prior written
13 * permission.
14 *
15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23 *
24 **************************************************************************************************/
29 #pragma once
30
31 #include <cublas_v2.h>
32 #include <cuda_fp16.h>
33 #include <stdint.h>
34
35 #include "cutlass/numeric_types.h"
36 #include "cutlass/complex.h"
37
38 namespace cutlass {
39 struct half_t;
40
41 template <typename T>
42 struct TypeTraits {
43 typedef T host_type;
44 typedef T device_type;
45 static inline T remove_negative_zero(T x) { return x; }
46 static inline T to_print(T x) { return x; }
47 static inline device_type to_device(host_type x) { return x; }
48 };
49
50 template <>
51 struct TypeTraits<int8_t> {
52 static cudaDataType_t const cublas_type = CUDA_R_8I;
53 typedef int8_t host_type;
54 typedef int8_t device_type;
55 typedef int8_t integer_type;
56 typedef uint8_t unsigned_type;
57 static inline int8_t remove_negative_zero(int8_t x) { return x; }
58 static inline int to_print(int8_t x) { return (int)x; }
59 static inline device_type to_device(host_type x) { return x; }
60 };
61
62 template <>
63 struct TypeTraits<uint8_t> {
64 static cudaDataType_t const cublas_type = CUDA_R_8I;
65 typedef uint8_t host_type;
66 typedef uint8_t device_type;
67 typedef uint8_t integer_type;
68 typedef uint8_t unsigned_type;
69 static inline uint8_t remove_negative_zero(uint8_t x) { return x; }
70 static inline uint32_t to_print(uint8_t x) { return (uint32_t)x; }
71 static inline device_type to_device(host_type x) { return x; }
72 };
73
74 template <>
75 struct TypeTraits<int> {
76 static cudaDataType_t const cublas_type = CUDA_R_32I;
77 typedef int host_type;
78 typedef int device_type;
79 typedef int32_t integer_type;
80 typedef uint32_t unsigned_type;
81 static inline int32_t remove_negative_zero(int32_t x) { return x; }
82 static inline int to_print(int x) { return x; }
83 static inline device_type to_device(host_type x) { return x; }
84 };
85
86 template <>
87 struct TypeTraits<unsigned> {
88 static cudaDataType_t const cublas_type = CUDA_R_32I;
89 typedef unsigned host_type;
90 typedef unsigned device_type;
91 typedef uint32_t integer_type;
92 typedef uint32_t unsigned_type;
93 static inline uint32_t remove_negative_zero(uint32_t x) { return x; }
94 static inline uint32_t to_print(uint32_t x) { return x; }
95 static inline device_type to_device(host_type x) { return x; }
96 };
97
98 template <>
99 struct TypeTraits<int64_t> {
100 static cudaDataType_t const cublas_type = CUDA_R_8I;
101 typedef int64_t host_type;
102 typedef int64_t device_type;
103 typedef int64_t integer_type;
104 typedef uint64_t unsigned_type;
105 static inline int64_t remove_negative_zero(int64_t x) { return x; }
106 static inline int64_t to_print(int64_t x) { return x; }
107 static inline device_type to_device(host_type x) { return x; }
108 };
109
110 template <>
111 struct TypeTraits<uint64_t> {
112 static cudaDataType_t const cublas_type = CUDA_R_8I;
113 typedef uint64_t host_type;
114 typedef uint64_t device_type;
115 typedef uint64_t integer_type;
116 typedef uint64_t unsigned_type;
117 static inline uint64_t remove_negative_zero(uint64_t x) { return x; }
118 static inline uint64_t to_print(uint64_t x) { return x; }
119 static inline device_type to_device(host_type x) { return x; }
120 };
121
122 template <>
123 struct TypeTraits<half_t> {
124 static cudaDataType_t const cublas_type = CUDA_R_16F;
125 typedef half_t host_type;
126 typedef half_t device_type;
127 typedef int16_t integer_type;
128 typedef uint16_t unsigned_type;
129 static inline half_t remove_negative_zero(half_t x) {
130 return (x.raw() == 0x8000 ? half_t::bitcast(0) : x);
131 }
132 static inline half_t to_print(half_t x) { return x; }
133 static inline device_type to_device(half_t x) { return reinterpret_cast<device_type const &>(x); }
134 };
135
136 template <>
137 struct TypeTraits<float> {
138 static cudaDataType_t const cublas_type = CUDA_R_32F;
139 typedef float host_type;
140 typedef float device_type;
141 typedef int32_t integer_type;
142 typedef uint32_t unsigned_type;
143 static inline float remove_negative_zero(float x) { return x == -0.f ? 0.f : x; }
144 static inline float to_print(float x) { return x; }
145 static inline device_type to_device(host_type x) { return x; }
146 };
147
148 template <>
149 struct TypeTraits<double> {
150 static cudaDataType_t const cublas_type = CUDA_R_64F;
151 typedef double host_type;
152 typedef double device_type;
153 typedef int64_t integer_type;
154 typedef uint64_t unsigned_type;
155 static inline double remove_negative_zero(double x) { return x == -0.0 ? 0.0 : x; }
156 static inline double to_print(double x) { return x; }
157 static inline device_type to_device(host_type x) { return x; }
158 };
159
161 //
162 // Complex types
163 //
165
166 template <>
167 struct TypeTraits<complex<half> > {
168 static cudaDataType_t const cublas_type = CUDA_C_16F;
169 typedef complex<half_t> host_type;
170 typedef complex<half> device_type;
171 typedef int16_t integer_type;
172 typedef uint16_t unsigned_type;
173 static inline device_type to_device(complex<half> x) { return reinterpret_cast<device_type const &>(x); }
174 };
175
176 template <>
177 struct TypeTraits<complex<half_t> > {
178 static cudaDataType_t const cublas_type = CUDA_C_16F;
179 typedef complex<half_t> host_type;
180 typedef complex<half> device_type;
181 typedef int16_t integer_type;
182 typedef uint16_t unsigned_type;
183 static inline complex<half_t> remove_negative_zero(complex<half_t> x) {
184 return complex<half_t>(
185 real(x) == -0_hf ? 0_hf : real(x),
186 imag(x) == -0_hf ? 0_hf : imag(x)
187 );
188 }
189 static inline complex<half_t> to_print(complex<half_t> x) { return x; }
190 static inline device_type to_device(complex<half_t> x) { return reinterpret_cast<device_type const &>(x); }
191 };
192
193 template <>
194 struct TypeTraits<complex<float> > {
195
196 static cudaDataType_t const cublas_type = CUDA_C_32F;
197 typedef complex<float> host_type;
198 typedef complex<float> device_type;
199 typedef int64_t integer_type;
200 typedef uint64_t unsigned_type;
201
202 static inline complex<float> remove_negative_zero(complex<float> x) {
203 return complex<float>(
204 real(x) == -0.f ? 0.f : real(x),
205 imag(x) == -0.f ? 0.f : imag(x)
206 );
207 }
208
209 static inline complex<float> to_print(complex<float> x) { return x; }
210 static inline device_type to_device(complex<float> x) { return reinterpret_cast<device_type const &>(x); }
211 };
212
213 template <>
214 struct TypeTraits<complex<double> > {
215 static cudaDataType_t const cublas_type = CUDA_C_64F;
216 typedef complex<double> host_type;
217 typedef complex<double> device_type;
218 struct integer_type { int64_t real, imag; };
219 struct unsigned_type { uint64_t real, imag; };
220 static inline complex<double> remove_negative_zero(complex<double> x) {
221 return complex<double>(
222 real(x) == -0.0 ? 0.0 : real(x),
223 imag(x) == -0.0 ? 0.0 : imag(x)
224 );
225 }
226 static inline complex<double> to_print(complex<double> x) { return x; }
227 static inline device_type to_device(complex<double> x) { return reinterpret_cast<device_type const &>(x); }
228 };
229
231
232 } // namespace cutlass
cutlass::TypeTraits< int64_t >::integer_type
int64_t integer_type
Definition: type_traits.h:103
cutlass::TypeTraits< complex< half > >::host_type
complex< half_t > host_type
Definition: type_traits.h:169
cutlass::half_t::bitcast
static CUTLASS_HOST_DEVICE half_t bitcast(uint16_t x)
Constructs from an unsigned short.
Definition: half.h:141
cutlass::TypeTraits< complex< float > >::to_print
static complex< float > to_print(complex< float > x)
Definition: type_traits.h:209
cutlass
Definition: aligned_buffer.h:35
cutlass::TypeTraits< double >::remove_negative_zero
static double remove_negative_zero(double x)
Definition: type_traits.h:155
cutlass::TypeTraits::host_type
T host_type
Definition: type_traits.h:43
cutlass::TypeTraits< half_t >::integer_type
int16_t integer_type
Definition: type_traits.h:127
cutlass::TypeTraits< double >::unsigned_type
uint64_t unsigned_type
Definition: type_traits.h:154
complex.h
cutlass::TypeTraits< float >::device_type
float device_type
Definition: type_traits.h:140
cutlass::imag
CUTLASS_HOST_DEVICE float const & imag(cuFloatComplex const &z)
Returns the imaginary part of the complex number.
Definition: complex.h:72
cutlass::TypeTraits< int8_t >::remove_negative_zero
static int8_t remove_negative_zero(int8_t x)
Definition: type_traits.h:57
cutlass::TypeTraits< float >::remove_negative_zero
static float remove_negative_zero(float x)
Definition: type_traits.h:143
cutlass::TypeTraits< unsigned >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:95
cutlass::TypeTraits< unsigned >::integer_type
uint32_t integer_type
Definition: type_traits.h:91
cutlass::TypeTraits< complex< half_t > >::integer_type
int16_t integer_type
Definition: type_traits.h:181
cutlass::TypeTraits::device_type
T device_type
Definition: type_traits.h:44
cutlass::TypeTraits< int >::integer_type
int32_t integer_type
Definition: type_traits.h:79
cutlass::TypeTraits< unsigned >::unsigned_type
uint32_t unsigned_type
Definition: type_traits.h:92
cutlass::TypeTraits< int >::device_type
int device_type
Definition: type_traits.h:78
cutlass::TypeTraits< int >::remove_negative_zero
static int32_t remove_negative_zero(int32_t x)
Definition: type_traits.h:81
cutlass::TypeTraits< unsigned >::remove_negative_zero
static uint32_t remove_negative_zero(uint32_t x)
Definition: type_traits.h:93
cutlass::TypeTraits< uint64_t >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:119
cutlass::TypeTraits< int64_t >::unsigned_type
uint64_t unsigned_type
Definition: type_traits.h:104
cutlass::half_t
IEEE half-precision floating-point type.
Definition: half.h:126
cutlass::TypeTraits< complex< double > >::device_type
complex< double > device_type
Definition: type_traits.h:217
cutlass::TypeTraits< complex< half > >::integer_type
int16_t integer_type
Definition: type_traits.h:171
cutlass::TypeTraits< uint64_t >::integer_type
uint64_t integer_type
Definition: type_traits.h:115
cutlass::TypeTraits< complex< double > >::to_print
static complex< double > to_print(complex< double > x)
Definition: type_traits.h:226
cutlass::real
CUTLASS_HOST_DEVICE float const & real(cuFloatComplex const &z)
Returns the real part of the complex number.
Definition: complex.h:56
cutlass::TypeTraits< half_t >::remove_negative_zero
static half_t remove_negative_zero(half_t x)
Definition: type_traits.h:129
cutlass::TypeTraits< complex< half_t > >::device_type
complex< half > device_type
Definition: type_traits.h:180
cutlass::TypeTraits< uint64_t >::host_type
uint64_t host_type
Definition: type_traits.h:113
cutlass::TypeTraits< float >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:145
cutlass::TypeTraits< int64_t >::host_type
int64_t host_type
Definition: type_traits.h:101
cutlass::TypeTraits< half_t >::to_print
static half_t to_print(half_t x)
Definition: type_traits.h:132
cutlass::TypeTraits< uint8_t >::integer_type
uint8_t integer_type
Definition: type_traits.h:67
cutlass::TypeTraits::remove_negative_zero
static T remove_negative_zero(T x)
Definition: type_traits.h:45
cutlass::TypeTraits< double >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:157
cutlass::TypeTraits< complex< half_t > >::host_type
complex< half_t > host_type
Definition: type_traits.h:179
cutlass::TypeTraits< float >::to_print
static float to_print(float x)
Definition: type_traits.h:144
cutlass::TypeTraits< int64_t >::device_type
int64_t device_type
Definition: type_traits.h:102
cutlass::TypeTraits< complex< double > >::remove_negative_zero
static complex< double > remove_negative_zero(complex< double > x)
Definition: type_traits.h:220
cutlass::TypeTraits< uint8_t >::remove_negative_zero
static uint8_t remove_negative_zero(uint8_t x)
Definition: type_traits.h:69
cutlass::TypeTraits< double >::device_type
double device_type
Definition: type_traits.h:152
cutlass::TypeTraits< int64_t >::remove_negative_zero
static int64_t remove_negative_zero(int64_t x)
Definition: type_traits.h:105
cutlass::TypeTraits< uint64_t >::to_print
static uint64_t to_print(uint64_t x)
Definition: type_traits.h:118
cutlass::TypeTraits< complex< float > >::integer_type
int64_t integer_type
Definition: type_traits.h:199
cutlass::TypeTraits< half_t >::to_device
static device_type to_device(half_t x)
Definition: type_traits.h:133
cutlass::TypeTraits< complex< double > >::host_type
complex< double > host_type
Definition: type_traits.h:216
cutlass::TypeTraits< uint64_t >::remove_negative_zero
static uint64_t remove_negative_zero(uint64_t x)
Definition: type_traits.h:117
cutlass::TypeTraits< float >::unsigned_type
uint32_t unsigned_type
Definition: type_traits.h:142
cutlass::TypeTraits< unsigned >::to_print
static uint32_t to_print(uint32_t x)
Definition: type_traits.h:94
cutlass::TypeTraits< float >::integer_type
int32_t integer_type
Definition: type_traits.h:141
cutlass::TypeTraits< int >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:83
cutlass::TypeTraits< int64_t >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:107
cutlass::TypeTraits< double >::host_type
double host_type
Definition: type_traits.h:151
cutlass::TypeTraits< int8_t >::device_type
int8_t device_type
Definition: type_traits.h:54
cutlass::TypeTraits< int >::unsigned_type
uint32_t unsigned_type
Definition: type_traits.h:80
cutlass::TypeTraits< complex< double > >::unsigned_type::real
uint64_t real
Definition: type_traits.h:219
cutlass::TypeTraits< complex< float > >::to_device
static device_type to_device(complex< float > x)
Definition: type_traits.h:210
cutlass::TypeTraits< unsigned >::device_type
unsigned device_type
Definition: type_traits.h:90
cutlass::TypeTraits< uint64_t >::device_type
uint64_t device_type
Definition: type_traits.h:114
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::TypeTraits< int64_t >::to_print
static int64_t to_print(int64_t x)
Definition: type_traits.h:106
cutlass::TypeTraits< int8_t >::unsigned_type
uint8_t unsigned_type
Definition: type_traits.h:56
cutlass::TypeTraits< uint8_t >::unsigned_type
uint8_t unsigned_type
Definition: type_traits.h:68
cutlass::TypeTraits::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:47
cutlass::TypeTraits< int8_t >::integer_type
int8_t integer_type
Definition: type_traits.h:55
cutlass::TypeTraits< int >::to_print
static int to_print(int x)
Definition: type_traits.h:82
cutlass::TypeTraits< complex< double > >::integer_type::real
int64_t real
Definition: type_traits.h:218
cutlass::TypeTraits< half_t >::unsigned_type
uint16_t unsigned_type
Definition: type_traits.h:128
cutlass::TypeTraits< uint8_t >::to_print
static uint32_t to_print(uint8_t x)
Definition: type_traits.h:70
cutlass::TypeTraits< double >::to_print
static double to_print(double x)
Definition: type_traits.h:156
cutlass::TypeTraits< complex< half_t > >::to_device
static device_type to_device(complex< half_t > x)
Definition: type_traits.h:190
cutlass::TypeTraits< unsigned >::host_type
unsigned host_type
Definition: type_traits.h:89
cutlass::TypeTraits< half_t >::host_type
half_t host_type
Definition: type_traits.h:125
cutlass::TypeTraits< uint64_t >::unsigned_type
uint64_t unsigned_type
Definition: type_traits.h:116
cutlass::TypeTraits< uint8_t >::host_type
uint8_t host_type
Definition: type_traits.h:65
cutlass::TypeTraits< double >::integer_type
int64_t integer_type
Definition: type_traits.h:153
cutlass::TypeTraits< complex< half > >::unsigned_type
uint16_t unsigned_type
Definition: type_traits.h:172
cutlass::complex
Definition: complex.h:92
cutlass::TypeTraits< complex< float > >::remove_negative_zero
static complex< float > remove_negative_zero(complex< float > x)
Definition: type_traits.h:202
cutlass::TypeTraits< float >::host_type
float host_type
Definition: type_traits.h:139
cutlass::TypeTraits< int8_t >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:59
cutlass::TypeTraits< complex< float > >::device_type
complex< float > device_type
Definition: type_traits.h:198
cutlass::TypeTraits< complex< float > >::host_type
complex< float > host_type
Definition: type_traits.h:197
cutlass::TypeTraits< complex< float > >::unsigned_type
uint64_t unsigned_type
Definition: type_traits.h:200
cutlass::TypeTraits< complex< double > >::to_device
static device_type to_device(complex< double > x)
Definition: type_traits.h:227
cutlass::TypeTraits< int >::host_type
int host_type
Definition: type_traits.h:77
cutlass::TypeTraits< uint8_t >::to_device
static device_type to_device(host_type x)
Definition: type_traits.h:71
cutlass::TypeTraits< complex< half_t > >::to_print
static complex< half_t > to_print(complex< half_t > x)
Definition: type_traits.h:189
cutlass::TypeTraits< complex< half_t > >::unsigned_type
uint16_t unsigned_type
Definition: type_traits.h:182
cutlass::TypeTraits< half_t >::device_type
half_t device_type
Definition: type_traits.h:126
cutlass::TypeTraits< int8_t >::to_print
static int to_print(int8_t x)
Definition: type_traits.h:58
cutlass::TypeTraits< complex< half > >::to_device
static device_type to_device(complex< half > x)
Definition: type_traits.h:173
cutlass::TypeTraits< complex< half > >::device_type
complex< half > device_type
Definition: type_traits.h:170
cutlass::TypeTraits< uint8_t >::device_type
uint8_t device_type
Definition: type_traits.h:66
cutlass::TypeTraits< int8_t >::host_type
int8_t host_type
Definition: type_traits.h:53
cutlass::TypeTraits< complex< half_t > >::remove_negative_zero
static complex< half_t > remove_negative_zero(complex< half_t > x)
Definition: type_traits.h:183
cutlass::half_t::raw
CUTLASS_HOST_DEVICE uint16_t & raw()
Accesses raw internal state.
Definition: half.h:367
cutlass::TypeTraits::to_print
static T to_print(T x)
Definition: type_traits.h:46
cutlass::TypeTraits
Definition: type_traits.h:42