30 #if !defined(__CUDACC_RTC__) || defined(CUTLASS_NVRTC_HAS_FP16) 31 #include <cuda_fp16.h> 41 template <
size_t kAlignment_>
61 template <
typename Scalar_,
int kLanes_>
121 #if !defined(__CUDACC_RTC__) || defined(CUTLASS_NVRTC_HAS_FP16) 123 template <
int kLanes_>
160 template <
int kLanes_>
173 "May only construct vectors of bin1_t that are multiples of 8 bits.");
187 return ( (
registers[i / 32] & (1 << (i % 32))) != 0 );
194 template <
int kLanes_>
207 "May only construct vectors of int4_t that are multiples of 8 bits.");
221 return (
registers[i / 8] >> (i % 8 * 4) & 0x0f)
222 - 16 * (
registers[i / 8] >> (i % 8 * 4 + 3) & 0x01);
229 template <
int kLanes_>
242 "May only construct vectors of uint4_t that are multiples of 8 bits.");
256 return registers[i / 8] >> (i % 8 * 4) & 0x0f;
262 template <
typename Scalar_>
269 template <
typename Element_,
int kLanes_ = 1>
276 template <
int kLanes_>
283 template <
int kLanes_>
290 template <
int kLanes_>
297 template <
typename Scalar_,
int kLanes_>
299 for (
int i = 0; i < Vector<Scalar_, kLanes_>::kRegisters; ++i) {
310 template <
typename T>
316 template <
typename T,
int Lanes>
322 template <
typename T,
int Lanes>
330 template <
typename T>
346 template <
typename T,
int Lanes>
362 template <
typename T,
int Lanes>
Vector< bin1_t, kLanes_ *32 > Type
Definition: vector.h:278
CUTLASS_HOST_DEVICE Scalar const & operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:84
CUTLASS_HOST_DEVICE void make_zero(Scalar_ &x)
Definition: vector.h:263
CUTLASS_HOST_DEVICE Vector(uint32_t value)
Constructor to convert from uint32_t type.
Definition: vector.h:218
half Scalar
The scalar type.
Definition: vector.h:94
Definition: numeric_types.h:39
T Scalar
Scalar type.
Definition: vector.h:333
CUTLASS_HOST_DEVICE int operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:220
CUTLASS_HOST_DEVICE Scalar const & operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:112
struct __align__(1) AlignedStruct< 1 >
Definition: vector.h:45
CUTLASS_HOST_DEVICE Scalar & operator[](uint32_t i)
Accessor to the ith lane.
Definition: vector.h:86
CUTLASS_HOST_DEVICE Vector(uint32_t value)
Constructor to convert from uint32_t type.
Definition: vector.h:253
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:242
Scalar_ Scalar
The scalar type.
Definition: vector.h:64
CUTLASS_HOST_DEVICE Scalar & operator[](uint32_t i)
Accessor to the ith lane.
Definition: vector.h:116
half Scalar
The scalar type.
Definition: vector.h:126
Vector< int4_t, kLanes_ *8 > Type
Definition: vector.h:285
uint32_t registers[kRegisters]
The data in registers.
Definition: vector.h:81
CUTLASS_HOST_DEVICE bool operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:186
uint4_t Scalar
The scalar type.
Definition: vector.h:232
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:104
bin1_t Scalar
The scalar type.
Definition: vector.h:163
CUTLASS_HOST_DEVICE Vector()
Default Constructor.
Definition: vector.h:216
Vector< T, 1 > Vector
Type that is always a vector.
Definition: vector.h:342
CUTLASS_HOST_DEVICE Vector(uint32_t value)
Constructor to convert from uint32_t type.
Definition: vector.h:184
CUTLASS_HOST_DEVICE Scalar & operator[](uint32_t i)
Accessor to the ith lane.
Definition: vector.h:150
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
Traits describing properties of vectors and scalar-as-vectors.
Definition: vector.h:331
static bool const IsVector
True if the type is actually a cutlass::Vector, otherwise false.
Definition: vector.h:339
Scalar scalars[kLanes]
The associated array of scalars.
Definition: vector.h:79
CUTLASS_HOST_DEVICE Vector()
Default Constructor.
Definition: vector.h:182
CUTLASS_HOST_DEVICE int operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:255
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:207
Vector< T, Lanes > Vector
Type that is always a Vector.
Definition: vector.h:374
static int const kLanes
Number of lanes of vector.
Definition: vector.h:336
T Scalar
Scalar type.
Definition: vector.h:365
Vector< Element_, kLanes_ > Type
Definition: vector.h:271
T Scalar
Scalar type.
Definition: vector.h:349
Definition: numeric_types.h:43
int4_t Scalar
The scalar type.
Definition: vector.h:197
static size_t const kValue
Definition: vector.h:312
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:74
Vector< uint4_t, kLanes_ *8 > Type
Definition: vector.h:292
Definition: numeric_types.h:41
CUTLASS_HOST_DEVICE Scalar const & operator[](uint32_t i) const
Accessor to the ith lane.
Definition: vector.h:146
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:136
Vector< T, Lanes > Vector
Type that is always a Vector.
Definition: vector.h:358
CUTLASS_HOST_DEVICE Vector()
Default Constructor.
Definition: vector.h:251
Returns the extent of a scalar or vector.
Definition: vector.h:311
AlignedStruct< kVectorSize > aligned_
The aligned storage to make sure we have good alignment.
Definition: vector.h:173