docs/integer__subbyte_8h_source.html
| | CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers |
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
45static int const kBits = Bits;
46
48static bool const kSigned = Signed;
49
51using T = typename std::conditional<kSigned, int, unsigned>::type;
52
55
57static Storage const kMask = Storage((1 << kBits) - 1);
58
59//
60// Data members
61//
62
64
65//
66// Methods
67//
68
71integer_subbyte() { }
72
75integer_subbyte(int value)
76 : storage(reinterpret_cast<Storage const &>(value) & kMask) {}
77
79integer_subbyte(unsigned value)
80 : storage(reinterpret_cast<Storage const &>(value) & kMask) {}
81
84integer_subbyte(double value) {
86 storage = reinterpret_cast<Storage const &>(tmp) & kMask;
87 }
88
92if (kSigned) {
93// Sign extend
94if (storage & Storage(1 << (kBits - 1))) {
95return T(storage) | ~T(kMask);
96 }
97 }
98return T(storage);
99 }
100
103bool operator==(integer_subbyte const &rhs) const {
104return storage == rhs.storage;
105 }
106
109bool operator!=(integer_subbyte const &rhs) const {
110return storage != rhs.storage;
111 }
112
115bool operator<=(integer_subbyte const &rhs) const {
116if (kSigned) {
117if (storage & (1 << (kBits - 1))) {
118return !(rhs.storage < storage);
119 }
120 }
121return storage < rhs.storage;
122 }
123
126bool operator<(integer_subbyte const &rhs) const {
127if (kSigned) {
128if (storage & (1 << (kBits - 1))) {
129return !(rhs.storage <= storage);
130 }
131 }
132return storage < rhs.storage;
133 }
134
137bool operator>=(integer_subbyte const &rhs) const {
138return !(*this < rhs);
139 }
140
143bool operator>(integer_subbyte const &rhs) const {
144return !(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](structcutlass_1_1sizeof bits_3_01uint1b t_01_4.html) struct sizeof_bits<uint1b_t> {
[165](structcutlass_1_1sizeof bits_3_01uint1b t_01_4.html#aa4debf31c59ca797f2850e874f76acb3)static int const value = 1;
166 };
167
169 template <>
[170](structcutlass_1_1sizeof bits_3_01int4b t_01_4.html) struct sizeof_bits<int4b_t> {
[171](structcutlass_1_1sizeof bits_3_01int4b t_01_4.html#a7f7e854d5be514fedd37caf08b01c301)static int const value = 4;
172 };
173
175 template <>
[176](structcutlass_1_1sizeof bits_3_01uint4b t_01_4.html) struct sizeof_bits<uint4b_t> {
[177](structcutlass_1_1sizeof bits_3_01uint4b t_01_4.html#a4faf5412c4a37ccef78a42d70f9be720)static int const value = 4;
178 };
179
181
182 } // namespace cutlass
Definition: aligned_buffer.h:35
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
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
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
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
#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
Generated by 1.8.11