array.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/numeric_types.h"
namespace cutlass {
template <
typename T,
int N,
bool RegisterSized = sizeof_bits<T>::value >= 32
>
class Array;
template <typename T, int N, bool RegisterSized>
struct sizeof_bits<Array<T, N, RegisterSized> > {
static int const value =
sizeof(typename Array<T, N, RegisterSized>::Storage) * 8 * Array<T, N, RegisterSized>::kStorageElements;
};
CUTLASS_HOST_DEVICE
constexpr bool ispow2(unsigned x) {
return !(x & (x - 1));
}
CUTLASS_HOST_DEVICE
constexpr unsigned floor_pow_2(unsigned x) {
return ispow2(x) ? x : floor_pow_2(x >> 1);
}
template <
typename T,
int N
>
class Array<T, N, true> {
public:
using Storage = T;
using Element = T;
//static std::size_t const kStorageElements = N;
static size_t const kStorageElements = N;
static size_t const kElements = N;
//
// C++ standard members
//
typedef T value_type;
typedef size_t size_type;
typedef ptrdiff_t difference_type;
typedef value_type &reference;
typedef value_type const & const_reference;
typedef value_type *pointer;
typedef value_type const * const_pointer;
//
// Iterators
//
class iterator {
T *ptr_;
public:
CUTLASS_HOST_DEVICE
iterator(): ptr_(nullptr) { }
CUTLASS_HOST_DEVICE
iterator(T *_ptr): ptr_(_ptr) { }
CUTLASS_HOST_DEVICE
iterator &operator++() {
++ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
iterator &operator--() {
--ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
iterator operator++(int) {
iterator ret(*this);
++ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
iterator operator--(int) {
iterator ret(*this);
--ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
T &operator*() const {
return *ptr_;
}
CUTLASS_HOST_DEVICE
bool operator==(iterator const &other) const {
return ptr_ == other.ptr_;
}
CUTLASS_HOST_DEVICE
bool operator!=(iterator const &other) const {
return ptr_ != other.ptr_;
}
};
class const_iterator {
T *ptr_;
public:
CUTLASS_HOST_DEVICE
const_iterator(): ptr_(nullptr) { }
CUTLASS_HOST_DEVICE
const_iterator(T const *_ptr): ptr_(_ptr) { }
CUTLASS_HOST_DEVICE
const_iterator &operator++() {
++ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
const_iterator &operator--() {
--ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
const_iterator operator++(int) {
const_iterator ret(*this);
++ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
const_iterator operator--(int) {
const_iterator ret(*this);
--ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
T const &operator*() const {
return *ptr_;
}
CUTLASS_HOST_DEVICE
bool operator==(const_iterator const &other) const {
return ptr_ == other.ptr_;
}
CUTLASS_HOST_DEVICE
bool operator!=(const_iterator const &other) const {
return ptr_ != other.ptr_;
}
};
class reverse_iterator {
T *ptr_;
public:
CUTLASS_HOST_DEVICE
reverse_iterator(): ptr_(nullptr) { }
CUTLASS_HOST_DEVICE
reverse_iterator(T *_ptr): ptr_(_ptr) { }
CUTLASS_HOST_DEVICE
reverse_iterator &operator++() {
--ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
reverse_iterator &operator--() {
++ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
reverse_iterator operator++(int) {
iterator ret(*this);
--ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
reverse_iterator operator--(int) {
iterator ret(*this);
++ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
T &operator*() const {
return *(ptr_ - 1);
}
CUTLASS_HOST_DEVICE
bool operator==(reverse_iterator const &other) const {
return ptr_ == other.ptr_;
}
CUTLASS_HOST_DEVICE
bool operator!=(reverse_iterator const &other) const {
return ptr_ != other.ptr_;
}
};
class const_reverse_iterator {
T const *ptr_;
public:
CUTLASS_HOST_DEVICE
const_reverse_iterator(): ptr_(nullptr) { }
CUTLASS_HOST_DEVICE
const_reverse_iterator(T const *_ptr): ptr_(_ptr) { }
CUTLASS_HOST_DEVICE
const_reverse_iterator &operator++() {
--ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
const_reverse_iterator &operator--() {
++ptr_;
return *this;
}
CUTLASS_HOST_DEVICE
const_reverse_iterator operator++(int) {
const_reverse_iterator ret(*this);
--ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
const_reverse_iterator operator--(int) {
const_reverse_iterator ret(*this);
++ptr_;
return ret;
}
CUTLASS_HOST_DEVICE
T const &operator*() const {
return *(ptr_ - 1);
}
CUTLASS_HOST_DEVICE
bool operator==(const_iterator const &other) const {
return ptr_ == other.ptr_;
}
CUTLASS_HOST_DEVICE
bool operator!=(const_iterator const &other) const {
return ptr_ != other.ptr_;
}
};
private:
Storage storage[kElements];
public:
CUTLASS_HOST_DEVICE
Array() { }
CUTLASS_HOST_DEVICE
Array(Array const &x) {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kElements; ++i) {
storage[i] = x.storage[i];
}
}
CUTLASS_HOST_DEVICE
void clear() {
fill(T(0));
}
CUTLASS_HOST_DEVICE
reference at(size_type pos) {
return reinterpret_cast<reference>(storage[pos]);
}
CUTLASS_HOST_DEVICE
const_reference at(size_type pos) const {
return reinterpret_cast<const_reference>(storage[pos]);
}
CUTLASS_HOST_DEVICE
reference operator[](size_type pos) {
return reinterpret_cast<reference>(storage[pos]);
}
CUTLASS_HOST_DEVICE
const_reference operator[](size_type pos) const {
return reinterpret_cast<const_reference>(storage[pos]);
}
CUTLASS_HOST_DEVICE
reference front() {
return reinterpret_cast<reference>(storage[0]);
}
CUTLASS_HOST_DEVICE
const_reference front() const {
return reinterpret_cast<const_reference>(storage[0]);
}
CUTLASS_HOST_DEVICE
reference back() {
return reinterpret_cast<reference>(storage[kStorageElements - 1]);
}
CUTLASS_HOST_DEVICE
const_reference back() const {
return reinterpret_cast<const_reference>(storage[kStorageElements - 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
pointer raw_data() {
return reinterpret_cast<pointer>(storage);
}
CUTLASS_HOST_DEVICE
const_pointer raw_data() const {
return reinterpret_cast<const_pointer>(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) {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kElements; ++i) {
storage[i] = static_cast< Storage >(value);
}
}
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(reinterpret_cast<pointer>(storage + kStorageElements));
}
CUTLASS_HOST_DEVICE
const_iterator cend() const {
return const_iterator(reinterpret_cast<const_pointer>(storage + kStorageElements));
}
CUTLASS_HOST_DEVICE
reverse_iterator rbegin() {
return reverse_iterator(reinterpret_cast<pointer>(storage + kStorageElements));
}
CUTLASS_HOST_DEVICE
const_reverse_iterator crbegin() const {
return const_reverse_iterator(reinterpret_cast<const_pointer>(storage + kStorageElements));
}
CUTLASS_HOST_DEVICE
reverse_iterator rend() {
return reverse_iterator(reinterpret_cast<pointer>(storage));
}
CUTLASS_HOST_DEVICE
const_reverse_iterator crend() const {
return const_reverse_iterator(reinterpret_cast<const_pointer>(storage));
}
//
// Comparison operators
//
};
} // namespace cutlass
#include "cutlass/array_subbyte.h"
namespace cutlass {
template <
typename T,
int N,
int Alignment = sizeof_bits<T>::value * N / 8
>
class alignas(Alignment) AlignedArray: public Array<T, N> {
public:
};
} // namespace cutlass
cutlass::Array< T, N, true >::pointer
value_type * pointer
Definition: array.h:103
cutlass::Array< T, N, true >::iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const
Definition: array.h:161
cutlass::Array< T, N, true >::const_reverse_iterator::const_reverse_iterator
CUTLASS_HOST_DEVICE const_reverse_iterator()
Definition: array.h:287
cutlass::Array< T, N, true >::const_iterator::const_iterator
CUTLASS_HOST_DEVICE const_iterator(T const *_ptr)
Definition: array.h:178
cutlass::Array< T, N, true >::const_iterator::const_iterator
CUTLASS_HOST_DEVICE const_iterator()
Definition: array.h:175
cutlass::Array< T, N, true >::empty
CUTLASS_HOST_DEVICE constexpr bool empty() const
Definition: array.h:420
cutlass::Array< T, N, true >::crbegin
CUTLASS_HOST_DEVICE const_reverse_iterator crbegin() const
Definition: array.h:468
cutlass::Array< T, N, true >::difference_type
ptrdiff_t difference_type
Definition: array.h:100
cutlass::Array< T, N, true >::reverse_iterator::operator--
CUTLASS_HOST_DEVICE reverse_iterator & operator--()
Definition: array.h:243
cutlass
Definition: aligned_buffer.h:35
constexpr
#define constexpr
Definition: platform.h:137
cutlass::sizeof_bits::value
static int const value
Definition: numeric_types.h:43
cutlass::Array< T, N, true >::const_iterator::operator++
CUTLASS_HOST_DEVICE const_iterator operator++(int)
Definition: array.h:193
cutlass::Array< T, N, true >::fill
CUTLASS_HOST_DEVICE void fill(T const &value)
Definition: array.h:435
array_subbyte.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
cutlass::Array< T, N, true >::end
CUTLASS_HOST_DEVICE iterator end()
Definition: array.h:453
cutlass::Array< T, N, true >::data
CUTLASS_HOST_DEVICE pointer data()
Definition: array.h:399
cutlass::Array< T, N, true >::begin
CUTLASS_HOST_DEVICE iterator begin()
Definition: array.h:443
cutlass::Array< T, N, true >::rbegin
CUTLASS_HOST_DEVICE reverse_iterator rbegin()
Definition: array.h:463
cutlass::AlignedArray
Aligned array type.
Definition: array.h:511
cutlass::Array< T, N, true >::operator[]
CUTLASS_HOST_DEVICE const_reference operator[](size_type pos) const
Definition: array.h:374
cutlass::Array< T, N, true >::raw_data
CUTLASS_HOST_DEVICE pointer raw_data()
Definition: array.h:409
cutlass::Array< T, N, true >::rend
CUTLASS_HOST_DEVICE reverse_iterator rend()
Definition: array.h:473
cutlass::Array< T, N, true >::cend
CUTLASS_HOST_DEVICE const_iterator cend() const
Definition: array.h:458
cutlass::Array< T, N, true >::const_reverse_iterator::const_reverse_iterator
CUTLASS_HOST_DEVICE const_reverse_iterator(T const *_ptr)
Definition: array.h:290
cutlass::Array< T, N, true >::reverse_iterator::operator--
CUTLASS_HOST_DEVICE reverse_iterator operator--(int)
Definition: array.h:256
cutlass::Array< T, N, true >::size_type
size_t size_type
Definition: array.h:99
cutlass::Array< T, N, true >::const_reverse_iterator::operator++
CUTLASS_HOST_DEVICE const_reverse_iterator & operator++()
Definition: array.h:293
cutlass::Array< T, N, true >::size
CUTLASS_HOST_DEVICE constexpr size_type size() const
Definition: array.h:425
cutlass::Array< T, N, true >::back
CUTLASS_HOST_DEVICE reference back()
Definition: array.h:389
cutlass::Array< T, N, true >::reverse_iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(reverse_iterator const &other) const
Definition: array.h:273
cutlass::Array< T, N, true >::reverse_iterator::operator++
CUTLASS_HOST_DEVICE reverse_iterator operator++(int)
Definition: array.h:249
cutlass::Array< T, N, true >::const_iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(const_iterator const &other) const
Definition: array.h:217
cutlass::Array< T, N, true >::const_reverse_iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(const_iterator const &other) const
Definition: array.h:329
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: cutlass.h:110
cutlass::Array< T, N, true >::const_reference
value_type const & const_reference
Definition: array.h:102
cutlass::Array< T, N, true >::reverse_iterator::operator++
CUTLASS_HOST_DEVICE reverse_iterator & operator++()
Definition: array.h:237
cutlass::Array< T, N, true >::const_iterator::operator*
CUTLASS_HOST_DEVICE T const & operator*() const
Definition: array.h:207
cutlass::Array< T, N, true >::reference
value_type & reference
Definition: array.h:101
cutlass::Array< T, N, true >::iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const
Definition: array.h:156
cutlass::Array< T, N, true >::crend
CUTLASS_HOST_DEVICE const_reverse_iterator crend() const
Definition: array.h:478
cutlass::Array< T, N, true >::reverse_iterator::reverse_iterator
CUTLASS_HOST_DEVICE reverse_iterator()
Definition: array.h:231
cutlass::Array< T, N, true >::iterator::operator*
CUTLASS_HOST_DEVICE T & operator*() const
Definition: array.h:151
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::Array< T, N, true >::iterator::operator++
CUTLASS_HOST_DEVICE iterator operator++(int)
Definition: array.h:137
cutlass::Array< T, N, true >::reverse_iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(reverse_iterator const &other) const
Definition: array.h:268
nullptr
#define nullptr
nullptr
Definition: platform.h:144
cutlass::Array< T, N, true >::iterator::operator++
CUTLASS_HOST_DEVICE iterator & operator++()
Definition: array.h:125
cutlass::Array< T, N, true >::Storage
T Storage
Storage type.
Definition: array.h:82
cutlass::Array< T, N, true >::reverse_iterator::reverse_iterator
CUTLASS_HOST_DEVICE reverse_iterator(T *_ptr)
Definition: array.h:234
cutlass::Array< T, N, true >::iterator::iterator
CUTLASS_HOST_DEVICE iterator(T *_ptr)
Definition: array.h:122
cutlass::Array< T, N, true >::const_iterator::operator++
CUTLASS_HOST_DEVICE const_iterator & operator++()
Definition: array.h:181
cutlass::Array< T, N, true >::back
CUTLASS_HOST_DEVICE const_reference back() const
Definition: array.h:394
cutlass::Array< T, N, true >::iterator::operator--
CUTLASS_HOST_DEVICE iterator & operator--()
Definition: array.h:131
cutlass::Array< T, N, true >::at
CUTLASS_HOST_DEVICE const_reference at(size_type pos) const
Definition: array.h:364
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::Array< T, N, true >::reverse_iterator::operator*
CUTLASS_HOST_DEVICE T & operator*() const
Definition: array.h:263
cutlass::Array< T, N, true >::const_iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(const_iterator const &other) const
Definition: array.h:212
cutlass::Array< T, N, true >::front
CUTLASS_HOST_DEVICE reference front()
Definition: array.h:379
cutlass::Array< T, N, true >::Element
T Element
Element type.
Definition: array.h:85
cutlass::Array< T, N, true >::raw_data
CUTLASS_HOST_DEVICE const_pointer raw_data() const
Definition: array.h:414
cutlass::Array< T, N, true >::max_size
CUTLASS_HOST_DEVICE constexpr size_type max_size() const
Definition: array.h:430
cutlass::Array< T, N, true >::data
CUTLASS_HOST_DEVICE const_pointer data() const
Definition: array.h:404
cutlass::Array< T, N, true >::const_reverse_iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(const_iterator const &other) const
Definition: array.h:324
cutlass::Array< T, N, true >::iterator::iterator
CUTLASS_HOST_DEVICE iterator()
Definition: array.h:119
cutlass::ispow2
CUTLASS_HOST_DEVICE constexpr bool ispow2(unsigned x)
Returns true if the argument is a power of 2.
Definition: array.h:59
cutlass::Array< T, N, true >::Array
CUTLASS_HOST_DEVICE Array(Array const &x)
Definition: array.h:345
cutlass::Array< T, N, true >::iterator::operator--
CUTLASS_HOST_DEVICE iterator operator--(int)
Definition: array.h:144
cutlass::floor_pow_2
CUTLASS_HOST_DEVICE constexpr unsigned floor_pow_2(unsigned x)
Returns the largest power of two not greater than the argument.
Definition: array.h:67
cutlass::Array< T, N, true >::const_reverse_iterator::operator--
CUTLASS_HOST_DEVICE const_reverse_iterator operator--(int)
Definition: array.h:312
cutlass::Array< T, N, true >::const_pointer
value_type const * const_pointer
Definition: array.h:104
cutlass::Array< T, N, true >::const_reverse_iterator::operator--
CUTLASS_HOST_DEVICE const_reverse_iterator & operator--()
Definition: array.h:299
cutlass::Array< T, N, true >::value_type
T value_type
Definition: array.h:98
cutlass::Array< T, N, true >::operator[]
CUTLASS_HOST_DEVICE reference operator[](size_type pos)
Definition: array.h:369
cutlass::Array< T, N, true >::cbegin
CUTLASS_HOST_DEVICE const_iterator cbegin() const
Definition: array.h:448
cutlass::Array< T, N, true >::Array
CUTLASS_HOST_DEVICE Array()
Definition: array.h:342
cutlass::Array< T, N, true >::const_iterator::operator--
CUTLASS_HOST_DEVICE const_iterator operator--(int)
Definition: array.h:200
cutlass::Array< T, N, true >::const_reverse_iterator::operator++
CUTLASS_HOST_DEVICE const_reverse_iterator operator++(int)
Definition: array.h:305
cutlass.h
Basic include for CUTLASS.
cutlass::Array< T, N, true >::const_reverse_iterator::operator*
CUTLASS_HOST_DEVICE T const & operator*() const
Definition: array.h:319
cutlass::Array< T, N, true >::front
CUTLASS_HOST_DEVICE const_reference front() const
Definition: array.h:384
cutlass::Array< T, N, true >::at
CUTLASS_HOST_DEVICE reference at(size_type pos)
Definition: array.h:359
cutlass::Array< T, N, true >::const_iterator::operator--
CUTLASS_HOST_DEVICE const_iterator & operator--()
Definition: array.h:187
cutlass::Array< T, N, true >::clear
CUTLASS_HOST_DEVICE void clear()
Efficient clear method.
Definition: array.h:354