array_subbyte.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 "cutlass/cutlass.h"
#include "cutlass/array.h"
#include "cutlass/platform/platform.h"
namespace cutlass {
template <
typename T,
int N
>
class Array<T, N, false> {
public:
static_assert(sizeof_bits<T>::value * N >= 8,
"Array<> specialized for sub-byte types assume the actual stored element size is 1 byte");
static int const kSizeBits = sizeof_bits<T>::value * N;
using Storage = typename platform::conditional<
((kSizeBits % 32) != 0),
typename platform::conditional<
((kSizeBits % 16) != 0),
uint8_t,
uint16_t
>::type,
uint32_t
>::type;
using Element = T;
static int const kElementsPerStoredItem = (sizeof(Storage) * 8) / sizeof_bits<T>::value;
static size_t const kStorageElements = N / kElementsPerStoredItem;
static size_t const kElements = N;
static Storage const kMask = ((Storage(1) << sizeof_bits<T>::value) - 1);
//
// C++ standard members with pointer types removed
//
typedef T value_type;
typedef size_t size_type;
typedef ptrdiff_t difference_type;
typedef value_type *pointer;
typedef value_type const *const_pointer;
//
// References
//
class reference {
Storage *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
reference(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
reference(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
CUTLASS_HOST_DEVICE
reference &operator=(T x) {
Storage item = (reinterpret_cast< Storage const &>(x) & kMask);
Storage kUpdateMask = Storage(~(kMask << (idx_ * sizeof_bits<T>::value)));
*ptr_ = Storage(((*ptr_ & kUpdateMask) | (item << idx_ * sizeof_bits<T>::value)));
return *this;
}
CUTLASS_HOST_DEVICE
T get() const {
Storage item = Storage((*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask);
return reinterpret_cast<T const &>(item);
}
CUTLASS_HOST_DEVICE
operator T() const {
return get();
}
CUTLASS_HOST_DEVICE
explicit operator int() const {
return int(get());
}
CUTLASS_HOST_DEVICE
explicit operator float() const {
return float(get());
}
};
class const_reference {
Storage const *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
const_reference(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
const_reference(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
CUTLASS_HOST_DEVICE
const T get() const {
Storage item = (*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask;
return reinterpret_cast<T const &>(item);
}
CUTLASS_HOST_DEVICE
operator T() const {
Storage item = Storage(Storage(*ptr_ >> Storage(idx_ * sizeof_bits<T>::value)) & kMask);
return reinterpret_cast<T const &>(item);
}
CUTLASS_HOST_DEVICE
explicit operator int() const {
return int(get());
}
CUTLASS_HOST_DEVICE
explicit operator float() const {
return float(get());
}
};
//
// Iterators
//
class iterator {
Storage *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
iterator(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
CUTLASS_HOST_DEVICE
iterator &operator++() {
++idx_;
if (idx_ == kElementsPerStoredItem) {
++ptr_;
idx_ = 0;
}
return *this;
}
CUTLASS_HOST_DEVICE
iterator &operator--() {
if (!idx_) {
--ptr_;
idx_ = kElementsPerStoredItem - 1;
}
else {
--idx_;
}
return *this;
}
CUTLASS_HOST_DEVICE
iterator operator++(int) {
iterator ret(*this);
++idx_;
if (idx_ == kElementsPerStoredItem) {
++ptr_;
idx_ = 0;
}
return ret;
}
CUTLASS_HOST_DEVICE
iterator operator--(int) {
iterator ret(*this);
if (!idx_) {
--ptr_;
idx_ = kElementsPerStoredItem - 1;
}
else {
--idx_;
}
return ret;
}
CUTLASS_HOST_DEVICE
reference operator*() const {
return reference(ptr_, idx_);
}
CUTLASS_HOST_DEVICE
bool operator==(iterator const &other) const {
return ptr_ == other.ptr_ && idx_ == other.idx_;
}
CUTLASS_HOST_DEVICE
bool operator!=(iterator const &other) const {
return !(*this == other);
}
};
class const_iterator {
Storage const *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
const_iterator(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
const_iterator(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
CUTLASS_HOST_DEVICE
iterator &operator++() {
++idx_;
if (idx_ == kElementsPerStoredItem) {
++ptr_;
idx_ = 0;
}
return *this;
}
CUTLASS_HOST_DEVICE
iterator &operator--() {
if (!idx_) {
--ptr_;
idx_ = kElementsPerStoredItem - 1;
}
else {
--idx_;
}
return *this;
}
CUTLASS_HOST_DEVICE
iterator operator++(int) {
iterator ret(*this);
++idx_;
if (idx_ == kElementsPerStoredItem) {
++ptr_;
idx_ = 0;
}
return ret;
}
CUTLASS_HOST_DEVICE
iterator operator--(int) {
iterator ret(*this);
if (!idx_) {
--ptr_;
idx_ = kElementsPerStoredItem - 1;
}
else {
--idx_;
}
return ret;
}
CUTLASS_HOST_DEVICE
const_reference operator*() const {
return const_reference(ptr_, idx_);
}
CUTLASS_HOST_DEVICE
bool operator==(iterator const &other) const {
return ptr_ == other.ptr_ && idx_ == other.idx_;
}
CUTLASS_HOST_DEVICE
bool operator!=(iterator const &other) const {
return !(*this == other);
}
};
class reverse_iterator {
Storage *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
reverse_iterator(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
reverse_iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
// TODO
};
class const_reverse_iterator {
Storage const *ptr_;
int idx_;
public:
CUTLASS_HOST_DEVICE
const_reverse_iterator(): ptr_(nullptr), idx_(0) { }
CUTLASS_HOST_DEVICE
const_reverse_iterator(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
// TODO
};
private:
Storage storage[kStorageElements];
public:
CUTLASS_HOST_DEVICE
Array() { }
CUTLASS_HOST_DEVICE
Array(Array const &x) {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < int(kStorageElements); ++i) {
storage[i] = x.storage[i];
}
}
CUTLASS_HOST_DEVICE
void clear() {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < int(kStorageElements); ++i) {
storage[i] = Storage(0);
}
}
CUTLASS_HOST_DEVICE
reference at(size_type pos) {
return reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);
}
CUTLASS_HOST_DEVICE
const_reference at(size_type pos) const {
return const_reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);
}
CUTLASS_HOST_DEVICE
reference operator[](size_type pos) {
return at(pos);
}
CUTLASS_HOST_DEVICE
const_reference operator[](size_type pos) const {
return at(pos);
}
CUTLASS_HOST_DEVICE
reference front() {
return at(0);
}
CUTLASS_HOST_DEVICE
const_reference front() const {
return at(0);
}
CUTLASS_HOST_DEVICE
reference back() {
return reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);
}
CUTLASS_HOST_DEVICE
const_reference back() const {
return const_reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);
}
CUTLASS_HOST_DEVICE
pointer data() {
return reinterpret_cast< pointer >(storage);
}
CUTLASS_HOST_DEVICE
const_pointer data() const {
return reinterpret_cast< const_pointer >(storage);
}
CUTLASS_HOST_DEVICE
Storage * raw_data() {
return storage;
}
CUTLASS_HOST_DEVICE
Storage const * raw_data() const {
return storage;
}
CUTLASS_HOST_DEVICE
constexpr bool empty() const {
return !kElements;
}
CUTLASS_HOST_DEVICE
constexpr size_type size() const {
return kElements;
}
CUTLASS_HOST_DEVICE
constexpr size_type max_size() const {
return kElements;
}
CUTLASS_HOST_DEVICE
void fill(T const &value) {
// TODO
}
CUTLASS_HOST_DEVICE
iterator begin() {
return iterator(storage);
}
CUTLASS_HOST_DEVICE
const_iterator cbegin() const {
return const_iterator(storage);
}
CUTLASS_HOST_DEVICE
iterator end() {
return iterator(storage + kStorageElements);
}
CUTLASS_HOST_DEVICE
const_iterator cend() const {
return const_iterator(storage + kStorageElements);
}
CUTLASS_HOST_DEVICE
reverse_iterator rbegin() {
return reverse_iterator(storage + kStorageElements);
}
CUTLASS_HOST_DEVICE
const_reverse_iterator crbegin() const {
return const_reverse_iterator(storage + kStorageElements);
}
CUTLASS_HOST_DEVICE
reverse_iterator rend() {
return reverse_iterator(storage);
}
CUTLASS_HOST_DEVICE
const_reverse_iterator crend() const {
return const_reverse_iterator(storage);
}
//
// Comparison operators
//
};
} // namespace cutlass
cutlass::Array< T, N, false >::const_reference::const_reference
CUTLASS_HOST_DEVICE const_reference(Storage const *ptr, int idx=0)
Ctor.
Definition: array_subbyte.h:164
cutlass::Array< T, N, false >::at
CUTLASS_HOST_DEVICE const_reference at(size_type pos) const
Definition: array_subbyte.h:428
cutlass::Array< T, N, false >::back
CUTLASS_HOST_DEVICE const_reference back() const
Definition: array_subbyte.h:458
cutlass
Definition: aligned_buffer.h:35
cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator
CUTLASS_HOST_DEVICE reverse_iterator()
Definition: array_subbyte.h:366
constexpr
#define constexpr
Definition: platform.h:137
cutlass::Array< T, N, false >::const_iterator::operator--
CUTLASS_HOST_DEVICE iterator operator--(int)
Definition: array_subbyte.h:326
cutlass::Array< T, N, false >::Array
CUTLASS_HOST_DEVICE Array(Array const &x)
Definition: array_subbyte.h:405
cutlass::Array< T, N, false >::operator[]
CUTLASS_HOST_DEVICE reference operator[](size_type pos)
Definition: array_subbyte.h:433
cutlass::Array< T, N, false >::crend
CUTLASS_HOST_DEVICE const_reverse_iterator crend() const
Definition: array_subbyte.h:539
cutlass::Array< T, N, false >::rbegin
CUTLASS_HOST_DEVICE reverse_iterator rbegin()
Definition: array_subbyte.h:524
cutlass::Array< T, N, false >::iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const
Definition: array_subbyte.h:265
cutlass::Array< T, N, false >::size_type
size_t size_type
Definition: array_subbyte.h:84
cutlass::Array< T, N, false >::iterator::operator++
CUTLASS_HOST_DEVICE iterator & operator++()
Definition: array_subbyte.h:214
cutlass::Array< T, N, false >::crbegin
CUTLASS_HOST_DEVICE const_reverse_iterator crbegin() const
Definition: array_subbyte.h:529
cutlass::Array< T, N, false >::rend
CUTLASS_HOST_DEVICE reverse_iterator rend()
Definition: array_subbyte.h:534
cutlass::Array< T, N, false >::empty
CUTLASS_HOST_DEVICE constexpr bool empty() const
Definition: array_subbyte.h:484
cutlass::Array< T, N, false >::const_iterator::const_iterator
CUTLASS_HOST_DEVICE const_iterator(Storage const *ptr, int idx=0)
Definition: array_subbyte.h:290
cutlass::Array< T, N, false >::difference_type
ptrdiff_t difference_type
Definition: array_subbyte.h:85
cutlass::Array< T, N, false >::data
CUTLASS_HOST_DEVICE const_pointer data() const
Definition: array_subbyte.h:468
cutlass::Array< T, N, false >::iterator::operator*
CUTLASS_HOST_DEVICE reference operator*() const
Definition: array_subbyte.h:260
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
cutlass::Array< T, N, false >::reference::reference
CUTLASS_HOST_DEVICE reference()
Default ctor.
Definition: array_subbyte.h:105
cutlass::Array< T, N, false >::cend
CUTLASS_HOST_DEVICE const_iterator cend() const
Definition: array_subbyte.h:519
cutlass::Array< T, N, false >::reference::reference
CUTLASS_HOST_DEVICE reference(Storage *ptr, int idx=0)
Ctor.
Definition: array_subbyte.h:109
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: cutlass.h:110
cutlass::Array< T, N, false >::raw_data
CUTLASS_HOST_DEVICE Storage * raw_data()
Definition: array_subbyte.h:473
cutlass::Array< T, N, false >::const_reference::const_reference
CUTLASS_HOST_DEVICE const_reference()
Default ctor.
Definition: array_subbyte.h:160
cutlass::Array< T, N, false >::back
CUTLASS_HOST_DEVICE reference back()
Definition: array_subbyte.h:453
cutlass::Array< T, N, false >::begin
CUTLASS_HOST_DEVICE iterator begin()
Definition: array_subbyte.h:504
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::Array< T, N, false >::fill
CUTLASS_HOST_DEVICE void fill(T const &value)
Definition: array_subbyte.h:499
nullptr
#define nullptr
nullptr
Definition: platform.h:144
cutlass::Array< T, N, false >::const_iterator::operator--
CUTLASS_HOST_DEVICE iterator & operator--()
Definition: array_subbyte.h:303
cutlass::Array< T, N, false >::Array
CUTLASS_HOST_DEVICE Array()
Definition: array_subbyte.h:402
cutlass::Array< T, N, false >::Element
T Element
Element type.
Definition: array_subbyte.h:65
cutlass::Array< T, N, false >::const_iterator::operator*
CUTLASS_HOST_DEVICE const_reference operator*() const
Definition: array_subbyte.h:339
cutlass::Array< T, N, false >::iterator::operator--
CUTLASS_HOST_DEVICE iterator operator--(int)
Definition: array_subbyte.h:247
cutlass::Array< T, N, false >::reference::operator=
CUTLASS_HOST_DEVICE reference & operator=(T x)
Assignment.
Definition: array_subbyte.h:113
cutlass::Array< T, N, false >::end
CUTLASS_HOST_DEVICE iterator end()
Definition: array_subbyte.h:514
cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator
CUTLASS_HOST_DEVICE reverse_iterator(Storage *ptr, int idx=0)
Definition: array_subbyte.h:369
cutlass::Array< T, N, false >::data
CUTLASS_HOST_DEVICE pointer data()
Definition: array_subbyte.h:463
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::platform::conditional
std::conditional (true specialization)
Definition: platform.h:325
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::Array< T, N, false >::front
CUTLASS_HOST_DEVICE reference front()
Definition: array_subbyte.h:443
cutlass::Array< T, N, false >::const_reverse_iterator::const_reverse_iterator
CUTLASS_HOST_DEVICE const_reverse_iterator(Storage const *ptr, int idx=0)
Definition: array_subbyte.h:389
cutlass::Array< T, N, false >::cbegin
CUTLASS_HOST_DEVICE const_iterator cbegin() const
Definition: array_subbyte.h:509
cutlass::Array< T, N, false >::const_iterator::const_iterator
CUTLASS_HOST_DEVICE const_iterator()
Definition: array_subbyte.h:287
cutlass::Array< T, N, false >::at
CUTLASS_HOST_DEVICE reference at(size_type pos)
Definition: array_subbyte.h:423
cutlass::Array< T, N, false >::const_reverse_iterator::const_reverse_iterator
CUTLASS_HOST_DEVICE const_reverse_iterator()
Definition: array_subbyte.h:386
cutlass::Array< T, N, false >::const_iterator::operator++
CUTLASS_HOST_DEVICE iterator operator++(int)
Definition: array_subbyte.h:315
cutlass::Array< T, N, false >::raw_data
CUTLASS_HOST_DEVICE Storage const * raw_data() const
Definition: array_subbyte.h:478
cutlass::Array< T, N, false >::const_pointer
value_type const * const_pointer
Definition: array_subbyte.h:87
cutlass::Array< T, N, false >::size
CUTLASS_HOST_DEVICE constexpr size_type size() const
Definition: array_subbyte.h:489
cutlass::Array< T, N, false >::max_size
CUTLASS_HOST_DEVICE constexpr size_type max_size() const
Definition: array_subbyte.h:494
cutlass::Array< T, N, false >::clear
CUTLASS_HOST_DEVICE void clear()
Efficient clear method.
Definition: array_subbyte.h:414
cutlass::Array< T, N, false >::front
CUTLASS_HOST_DEVICE const_reference front() const
Definition: array_subbyte.h:448
cutlass::Array< T, N, false >::iterator::iterator
CUTLASS_HOST_DEVICE iterator()
Definition: array_subbyte.h:208
cutlass::Array< T, N, false >::iterator::iterator
CUTLASS_HOST_DEVICE iterator(Storage *ptr, int idx=0)
Definition: array_subbyte.h:211
cutlass::Array< T, N, false >::Storage
typename platform::conditional< ((kSizeBits%32)!=0), typename platform::conditional< ((kSizeBits%16)!=0), uint8_t, uint16_t >::type, uint32_t >::type Storage
Storage type.
Definition: array_subbyte.h:62
cutlass::Array< T, N, false >::iterator::operator++
CUTLASS_HOST_DEVICE iterator operator++(int)
Definition: array_subbyte.h:236
cutlass::Array< T, N, false >::iterator::operator--
CUTLASS_HOST_DEVICE iterator & operator--()
Definition: array_subbyte.h:224
cutlass::Array< T, N, false >::const_iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const
Definition: array_subbyte.h:349
cutlass::Array< T, N, false >::iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const
Definition: array_subbyte.h:270
cutlass::Array< T, N, false >::value_type
T value_type
Definition: array_subbyte.h:83
cutlass::Array< T, N, false >::pointer
value_type * pointer
Definition: array_subbyte.h:86
cutlass::Array< T, N, false >::const_iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const
Definition: array_subbyte.h:344
cutlass::Array< T, N, false >::const_iterator::operator++
CUTLASS_HOST_DEVICE iterator & operator++()
Definition: array_subbyte.h:293
cutlass.h
Basic include for CUTLASS.
cutlass::Array< T, N, false >::operator[]
CUTLASS_HOST_DEVICE const_reference operator[](size_type pos) const
Definition: array_subbyte.h:438