type_traits.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-2019, 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
#include <cublas_v2.h>
#include <cuda_fp16.h>
#include <stdint.h>
#include "cutlass/numeric_types.h"
#include "cutlass/complex.h"
namespace cutlass {
struct half_t;
template <typename T>
struct TypeTraits {
typedef T host_type;
typedef T device_type;
static inline T remove_negative_zero(T x) { return x; }
static inline T to_print(T x) { return x; }
static inline device_type to_device(host_type x) { return x; }
};
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; }
static inline device_type to_device(host_type x) { return 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; }
static inline device_type to_device(host_type x) { return 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; }
static inline device_type to_device(host_type 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; }
static inline device_type to_device(host_type 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; }
static inline device_type to_device(host_type 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; }
static inline device_type to_device(host_type x) { return x; }
};
template <>
struct TypeTraits<half_t> {
static cudaDataType_t const cublas_type = CUDA_R_16F;
typedef half_t host_type;
typedef half_t 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; }
static inline device_type to_device(half_t x) { return reinterpret_cast<device_type const &>(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; }
static inline device_type to_device(host_type 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; }
static inline device_type to_device(host_type x) { return x; }
};
//
// Complex types
//
template <>
struct TypeTraits<complex<half> > {
static cudaDataType_t const cublas_type = CUDA_C_16F;
typedef complex<half_t> host_type;
typedef complex<half> device_type;
typedef int16_t integer_type;
typedef uint16_t unsigned_type;
static inline device_type to_device(complex<half> x) { return reinterpret_cast<device_type const &>(x); }
};
template <>
struct TypeTraits<complex<half_t> > {
static cudaDataType_t const cublas_type = CUDA_C_16F;
typedef complex<half_t> host_type;
typedef complex<half> device_type;
typedef int16_t integer_type;
typedef uint16_t unsigned_type;
static inline complex<half_t> remove_negative_zero(complex<half_t> x) {
return complex<half_t>(
real(x) == -0_hf ? 0_hf : real(x),
imag(x) == -0_hf ? 0_hf : imag(x)
);
}
static inline complex<half_t> to_print(complex<half_t> x) { return x; }
static inline device_type to_device(complex<half_t> x) { return reinterpret_cast<device_type const &>(x); }
};
template <>
struct TypeTraits<complex<float> > {
static cudaDataType_t const cublas_type = CUDA_C_32F;
typedef complex<float> host_type;
typedef complex<float> device_type;
typedef int64_t integer_type;
typedef uint64_t unsigned_type;
static inline complex<float> remove_negative_zero(complex<float> x) {
return complex<float>(
real(x) == -0.f ? 0.f : real(x),
imag(x) == -0.f ? 0.f : imag(x)
);
}
static inline complex<float> to_print(complex<float> x) { return x; }
static inline device_type to_device(complex<float> x) { return reinterpret_cast<device_type const &>(x); }
};
template <>
struct TypeTraits<complex<double> > {
static cudaDataType_t const cublas_type = CUDA_C_64F;
typedef complex<double> host_type;
typedef complex<double> device_type;
struct integer_type { int64_t real, imag; };
struct unsigned_type { uint64_t real, imag; };
static inline complex<double> remove_negative_zero(complex<double> x) {
return complex<double>(
real(x) == -0.0 ? 0.0 : real(x),
imag(x) == -0.0 ? 0.0 : imag(x)
);
}
static inline complex<double> to_print(complex<double> x) { return x; }
static inline device_type to_device(complex<double> x) { return reinterpret_cast<device_type const &>(x); }
};
} // 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