| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

integer_subbyte.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 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 **************************************************************************************************/
30 #pragma once
31
32 #include <cstdint>
33
34 #include "cutlass/platform/platform.h"
35
36 namespace cutlass {
37
39
41 template <int Bits, bool Signed = true>
42 struct integer_subbyte {
43
45 static int const kBits = Bits;
46
48 static bool const kSigned = Signed;
49
51 using T = typename std::conditional<kSigned, int, unsigned>::type;
52
54 using Storage = uint8_t;
55
57 static Storage const kMask = Storage((1 << kBits) - 1);
58
59 //
60 // Data members
61 //
62
63 Storage storage;
64
65 //
66 // Methods
67 //
68
70 CUTLASS_HOST_DEVICE
71 integer_subbyte() { }
72
74 CUTLASS_HOST_DEVICE
75 integer_subbyte(int value)
76 : storage(reinterpret_cast<Storage const &>(value) & kMask) {}
77
78 CUTLASS_HOST_DEVICE
79 integer_subbyte(unsigned value)
80 : storage(reinterpret_cast<Storage const &>(value) & kMask) {}
81
83 CUTLASS_HOST_DEVICE
84 integer_subbyte(double value) {
85 T tmp = (T)value;
86 storage = reinterpret_cast<Storage const &>(tmp) & kMask;
87 }
88
90 CUTLASS_HOST_DEVICE
91 operator T() const {
92 if (kSigned) {
93 // Sign extend
94 if (storage & Storage(1 << (kBits - 1))) {
95 return T(storage) | ~T(kMask);
96 }
97 }
98 return T(storage);
99 }
100
102 CUTLASS_HOST_DEVICE
103 bool operator==(integer_subbyte const &rhs) const {
104 return storage == rhs.storage;
105 }
106
108 CUTLASS_HOST_DEVICE
109 bool operator!=(integer_subbyte const &rhs) const {
110 return storage != rhs.storage;
111 }
112
114 CUTLASS_HOST_DEVICE
115 bool operator<=(integer_subbyte const &rhs) const {
116 if (kSigned) {
117 if (storage & (1 << (kBits - 1))) {
118 return !(rhs.storage < storage);
119 }
120 }
121 return storage < rhs.storage;
122 }
123
125 CUTLASS_HOST_DEVICE
126 bool operator<(integer_subbyte const &rhs) const {
127 if (kSigned) {
128 if (storage & (1 << (kBits - 1))) {
129 return !(rhs.storage <= storage);
130 }
131 }
132 return storage < rhs.storage;
133 }
134
136 CUTLASS_HOST_DEVICE
137 bool operator>=(integer_subbyte const &rhs) const {
138 return !(*this < rhs);
139 }
140
142 CUTLASS_HOST_DEVICE
143 bool operator>(integer_subbyte const &rhs) const {
144 return !(rhs < *this);
145 }
146 };
147
149
150
152 using uint1b_t = integer_subbyte<1, false>;
153
155 using int4b_t = integer_subbyte<4, true>;
156
158 using uint4b_t = integer_subbyte<4, false>;
159
161
163 template <>
164 struct sizeof_bits<uint1b_t> {
165 static int const value = 1;
166 };
167
169 template <>
170 struct sizeof_bits<int4b_t> {
171 static int const value = 4;
172 };
173
175 template <>
176 struct sizeof_bits<uint4b_t> {
177 static int const value = 4;
178 };
179
181
182 } // namespace cutlass
cutlass
Definition: aligned_buffer.h:35
cutlass::integer_subbyte
4-bit signed integer type
Definition: integer_subbyte.h:42
cutlass::integer_subbyte::integer_subbyte
CUTLASS_HOST_DEVICE integer_subbyte(int value)
Conversion from integer type.
Definition: integer_subbyte.h:75
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
cutlass::integer_subbyte::kMask
static Storage const kMask
Bitmask used to truncate from larger integers.
Definition: integer_subbyte.h:57
cutlass::integer_subbyte::operator==
CUTLASS_HOST_DEVICE bool operator==(integer_subbyte const &rhs) const
Equality.
Definition: integer_subbyte.h:103
cutlass::integer_subbyte::T
typename std::conditional< kSigned, int, unsigned >::type T
External type.
Definition: integer_subbyte.h:51
cutlass::integer_subbyte::kSigned
static bool const kSigned
Whether type is signed.
Definition: integer_subbyte.h:48
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::integer_subbyte::integer_subbyte
CUTLASS_HOST_DEVICE integer_subbyte(unsigned value)
Definition: integer_subbyte.h:79
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::integer_subbyte::kBits
static int const kBits
Number of bits.
Definition: integer_subbyte.h:45
cutlass::integer_subbyte::storage
Storage storage
Definition: integer_subbyte.h:63
cutlass::integer_subbyte::integer_subbyte
CUTLASS_HOST_DEVICE integer_subbyte()
No operation.
Definition: integer_subbyte.h:71
cutlass::integer_subbyte::operator>=
CUTLASS_HOST_DEVICE bool operator>=(integer_subbyte const &rhs) const
Greater than or equal.
Definition: integer_subbyte.h:137
cutlass::integer_subbyte::operator!=
CUTLASS_HOST_DEVICE bool operator!=(integer_subbyte const &rhs) const
Inequality.
Definition: integer_subbyte.h:109
cutlass::integer_subbyte::integer_subbyte
CUTLASS_HOST_DEVICE integer_subbyte(double value)
Conversion from double.
Definition: integer_subbyte.h:84
cutlass::integer_subbyte::Storage
uint8_t Storage
Storage type.
Definition: integer_subbyte.h:54
cutlass::integer_subbyte::operator<=
CUTLASS_HOST_DEVICE bool operator<=(integer_subbyte const &rhs) const
Less than or equal.
Definition: integer_subbyte.h:115
cutlass::integer_subbyte::operator>
CUTLASS_HOST_DEVICE bool operator>(integer_subbyte const &rhs) const
Greater than.
Definition: integer_subbyte.h:143
cutlass::integer_subbyte::operator<
CUTLASS_HOST_DEVICE bool operator<(integer_subbyte const &rhs) const
Less than.
Definition: integer_subbyte.h:126