integer_subbyte.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 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 <cstdint>
#include "cutlass/platform/platform.h"
namespace cutlass {
template <int Bits, bool Signed = true>
struct integer_subbyte {
static int const kBits = Bits;
static bool const kSigned = Signed;
using T = typename std::conditional<kSigned, int, unsigned>::type;
using Storage = uint8_t;
static Storage const kMask = Storage((1 << kBits) - 1);
//
// Data members
//
Storage storage;
//
// Methods
//
CUTLASS_HOST_DEVICE
integer_subbyte() { }
CUTLASS_HOST_DEVICE
integer_subbyte(int value)
: storage(reinterpret_cast<Storage const &>(value) & kMask) {}
CUTLASS_HOST_DEVICE
integer_subbyte(unsigned value)
: storage(reinterpret_cast<Storage const &>(value) & kMask) {}
CUTLASS_HOST_DEVICE
integer_subbyte(double value) {
T tmp = (T)value;
storage = reinterpret_cast< Storage const &>(tmp) & kMask;
}
CUTLASS_HOST_DEVICE
operator T() const {
if (kSigned) {
// Sign extend
if (storage & Storage(1 << (kBits - 1))) {
return T(storage) | ~T(kMask);
}
}
return T(storage);
}
CUTLASS_HOST_DEVICE
bool operator==(integer_subbyte const &rhs) const {
return storage == rhs.storage;
}
CUTLASS_HOST_DEVICE
bool operator!=(integer_subbyte const &rhs) const {
return storage != rhs.storage;
}
CUTLASS_HOST_DEVICE
bool operator<=(integer_subbyte const &rhs) const {
if (kSigned) {
if (storage & (1 << (kBits - 1))) {
return !(rhs.storage < storage);
}
}
return storage < rhs.storage;
}
CUTLASS_HOST_DEVICE
bool operator<(integer_subbyte const &rhs) const {
if (kSigned) {
if (storage & (1 << (kBits - 1))) {
return !(rhs.storage <= storage);
}
}
return storage < rhs.storage;
}
CUTLASS_HOST_DEVICE
bool operator>=(integer_subbyte const &rhs) const {
return !(*this < rhs);
}
CUTLASS_HOST_DEVICE
bool operator>(integer_subbyte const &rhs) const {
return !(rhs < *this);
}
};
using uint1b_t = integer_subbyte<1, false>;
using int4b_t = integer_subbyte<4, true>;
using uint4b_t = integer_subbyte<4, false>;
template <>
struct sizeof_bits<uint1b_t> {
static int const value = 1;
};
template <>
struct sizeof_bits<int4b_t> {
static int const value = 4;
};
template <>
struct sizeof_bits<uint4b_t> {
static int const value = 4;
};
} // 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