predicate_vector.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
#if !defined(__CUDACC_RTC__)
#include <assert.h>
#endif
#include <stdint.h>
#include "cutlass/cutlass.h"
#include "cutlass/platform/platform.h"
namespace cutlass {
template <
int kPredicates_,
int kPredicatesPerByte_ = 4,
int kPredicateStart_ = 0>
struct PredicateVector {
static int const kPredicates = kPredicates_;
static int const kPredicatesPerByte = kPredicatesPerByte_;
static int const kPredicateStart = kPredicateStart_;
// Make sure no one tries to put more than 8 bits in a byte :)
static_assert(kPredicatesPerByte <= 8, "kPredicatesPerByte must fit within an actual byte");
// Make sure the "offsetted" bits fit in one byte.
static_assert(kPredicateStart + kPredicatesPerByte <= 8,
"The offsetted predicates must fit within an actual byte.");
typedef uint32_t Storage;
static int const kBytes = (kPredicates + kPredicatesPerByte - 1) / kPredicatesPerByte;
static int const kWordCount = (kBytes + sizeof(Storage) - 1) / sizeof(Storage);
private:
//
// Data members
//
Storage storageData[kWordCount];
//
// Methods
//
CUTLASS_HOST_DEVICE void computeStorageOffset(int &word, int &bit, int idx) const {
CUTLASS_ASSERT(idx < kPredicates);
int byte = (idx / kPredicatesPerByte);
int bit_offset = (idx % kPredicatesPerByte);
word = byte / sizeof(Storage);
int byte_offset = (byte % sizeof(Storage));
bit = byte_offset * 8 + bit_offset + kPredicateStart;
}
CUTLASS_HOST_DEVICE Storage &storage(int word) {
CUTLASS_ASSERT(word < kWordCount);
return storageData[word];
}
CUTLASS_HOST_DEVICE Storage const &storage(int word) const {
CUTLASS_ASSERT(word < kWordCount);
return storageData[word];
}
public:
//
// Iterator
//
class Iterator {
PredicateVector &vec_;
int bit_;
public:
CUTLASS_HOST_DEVICE
Iterator(Iterator const &it) : vec_(it.vec_), bit_(it.bit_) {}
CUTLASS_HOST_DEVICE
Iterator(PredicateVector &vec, int _start = 0) : vec_(vec), bit_(_start) {}
CUTLASS_HOST_DEVICE
Iterator &operator++() {
++bit_;
return *this;
}
CUTLASS_HOST_DEVICE
Iterator &operator+=(int offset) {
bit_ += offset;
return *this;
}
CUTLASS_HOST_DEVICE
Iterator &operator--() {
--bit_;
return *this;
}
CUTLASS_HOST_DEVICE
Iterator &operator-=(int offset) {
bit_ -= offset;
return *this;
}
CUTLASS_HOST_DEVICE
Iterator operator++(int) {
Iterator ret(*this);
ret.bit_++;
return ret;
}
CUTLASS_HOST_DEVICE
Iterator operator--(int) {
Iterator ret(*this);
ret.bit_--;
return ret;
}
CUTLASS_HOST_DEVICE
Iterator operator+(int offset) {
Iterator ret(*this);
ret.bit_ += offset;
return ret;
}
CUTLASS_HOST_DEVICE
Iterator operator-(int offset) {
ConstIterator ret(*this);
ret.bit_ -= offset;
return ret;
}
CUTLASS_HOST_DEVICE
bool operator==(Iterator const &it) const { return bit_ == it.bit_; }
CUTLASS_HOST_DEVICE
bool operator!=(Iterator const &it) const { return bit_ != it.bit_; }
CUTLASS_HOST_DEVICE
bool get() { return vec_.at(bit_); }
CUTLASS_HOST_DEVICE
bool at() const { return vec_.at(bit_); }
CUTLASS_HOST_DEVICE
bool operator*() const { return at(); }
CUTLASS_HOST_DEVICE
void set(bool value = true) { vec_.set(bit_, value); }
};
class ConstIterator {
PredicateVector const &vec_;
int bit_;
public:
CUTLASS_HOST_DEVICE
ConstIterator(ConstIterator const &it) : vec_(it.vec_), bit_(it.bit_) {}
CUTLASS_HOST_DEVICE
ConstIterator(PredicateVector const &vec, int _start = 0) : vec_(vec), bit_(_start) {}
CUTLASS_HOST_DEVICE
ConstIterator &operator++() {
++bit_;
return *this;
}
CUTLASS_HOST_DEVICE
ConstIterator &operator+=(int offset) {
bit_ += offset;
return *this;
}
CUTLASS_HOST_DEVICE
ConstIterator &operator--() {
--bit_;
return *this;
}
CUTLASS_HOST_DEVICE
ConstIterator &operator-=(int offset) {
bit_ -= offset;
return *this;
}
CUTLASS_HOST_DEVICE
ConstIterator operator++(int) {
ConstIterator ret(*this);
ret.bit_++;
return ret;
}
CUTLASS_HOST_DEVICE
ConstIterator operator--(int) {
ConstIterator ret(*this);
ret.bit_--;
return ret;
}
CUTLASS_HOST_DEVICE
ConstIterator operator+(int offset) {
ConstIterator ret(*this);
ret.bit_ += offset;
return ret;
}
CUTLASS_HOST_DEVICE
ConstIterator operator-(int offset) {
ConstIterator ret(*this);
ret.bit_ -= offset;
return ret;
}
CUTLASS_HOST_DEVICE
bool operator==(ConstIterator const &it) const { return bit_ == it.bit_; }
CUTLASS_HOST_DEVICE
bool operator!=(ConstIterator const &it) const { return bit_ != it.bit_; }
CUTLASS_HOST_DEVICE
bool get() { return vec_.at(bit_); }
CUTLASS_HOST_DEVICE
bool at() const { return vec_.at(bit_); }
CUTLASS_HOST_DEVICE
bool operator*() const { return at(); }
};
struct TrivialIterator {
CUTLASS_HOST_DEVICE
TrivialIterator() {}
CUTLASS_HOST_DEVICE
TrivialIterator(Iterator const &it) {}
CUTLASS_HOST_DEVICE
TrivialIterator(PredicateVector const &_vec) {}
CUTLASS_HOST_DEVICE
TrivialIterator &operator++() { return *this; }
CUTLASS_HOST_DEVICE
TrivialIterator operator++(int) { return *this; }
CUTLASS_HOST_DEVICE
bool operator*() const { return true; }
};
public:
//
// Methods
//
CUTLASS_HOST_DEVICE PredicateVector(bool value = true) { fill(value); }
CUTLASS_HOST_DEVICE void fill(bool value = true) {
Storage item = (value ? ~Storage(0) : Storage(0));
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kWordCount; ++i) {
storage(i) = item;
}
}
CUTLASS_HOST_DEVICE void clear() {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kWordCount; ++i) {
storage(i) = 0;
}
}
CUTLASS_HOST_DEVICE void enable() {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kWordCount; ++i) {
storage(i) = ~Storage(0);
}
}
CUTLASS_HOST_DEVICE bool operator[](int idx) const { return at(idx); }
CUTLASS_HOST_DEVICE bool at(int idx) const {
int bit, word;
computeStorageOffset(word, bit, idx);
return ((storage(word) >> bit) & 1);
}
CUTLASS_HOST_DEVICE void set(int idx, bool value = true) {
int bit, word;
computeStorageOffset(word, bit, idx);
Storage disable_mask = (~(Storage(1) << bit));
Storage enable_mask = (Storage(value) << bit);
storage(word) = ((storage(word) & disable_mask) | enable_mask);
}
CUTLASS_HOST_DEVICE PredicateVector &operator&=(PredicateVector const &predicates) {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kWordCount; ++i) {
storage(i) = (storage(i) & predicates.storage(i));
}
return *this;
}
CUTLASS_HOST_DEVICE PredicateVector &operator|=(PredicateVector const &predicates) {
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < kWordCount; ++i) {
storage(i) = (storage(i) | predicates.storage(i));
}
return *this;
}
CUTLASS_HOST_DEVICE bool is_zero() const {
Storage mask(0);
for (int byte = 0; byte < sizeof(Storage); ++byte) {
Storage byte_mask = (((1 << kPredicatesPerByte) - 1) << kPredicateStart);
mask |= (byte_mask << (byte * 8));
}
uint32_t result = 0;
for (int word = 0; word < kWordCount; ++word) {
result |= storage(word);
}
return result == 0;
}
CUTLASS_DEVICE
Iterator begin() { return Iterator(*this); }
CUTLASS_DEVICE
Iterator end() { return Iterator(*this, kPredicates); }
CUTLASS_DEVICE
ConstIterator const_begin() const { return ConstIterator(*this); }
CUTLASS_DEVICE
ConstIterator const_end() const { return ConstIterator(*this, kPredicates); }
};
} // namespace cutlass
cutlass::PredicateVector::operator|=
CUTLASS_HOST_DEVICE PredicateVector & operator|=(PredicateVector const &predicates)
Computes the union of two identical predicate vectors.
Definition: predicate_vector.h:474
cutlass::PredicateVector::TrivialIterator::operator++
CUTLASS_HOST_DEVICE TrivialIterator & operator++()
Pre-increment.
Definition: predicate_vector.h:397
cutlass::PredicateVector::ConstIterator::ConstIterator
CUTLASS_HOST_DEVICE ConstIterator(PredicateVector const &vec, int _start=0)
Constructs an iterator from a PredicateVector.
Definition: predicate_vector.h:298
cutlass
Definition: aligned_buffer.h:35
cutlass::PredicateVector::Storage
uint32_t Storage
Storage type of individual elements.
Definition: predicate_vector.h:117
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator(PredicateVector const &_vec)
Constructs an iterator from a PredicateVector.
Definition: predicate_vector.h:393
cutlass::PredicateVector::is_zero
CUTLASS_HOST_DEVICE bool is_zero() const
Returns true if entire predicate array is zero.
Definition: predicate_vector.h:483
cutlass::PredicateVector::ConstIterator::operator--
CUTLASS_HOST_DEVICE ConstIterator & operator--()
Pre-decrement.
Definition: predicate_vector.h:316
cutlass::PredicateVector::kBytes
static int const kBytes
Number of bytes needed.
Definition: predicate_vector.h:126
cutlass::PredicateVector::const_end
CUTLASS_DEVICE ConstIterator const_end() const
Returns a ConstIterator.
Definition: predicate_vector.h:510
cutlass::PredicateVector::ConstIterator::at
CUTLASS_HOST_DEVICE bool at() const
Gets the bit at the pointed to location.
Definition: predicate_vector.h:374
cutlass::PredicateVector::ConstIterator::operator++
CUTLASS_HOST_DEVICE ConstIterator & operator++()
Pre-increment.
Definition: predicate_vector.h:302
cutlass::PredicateVector::enable
CUTLASS_HOST_DEVICE void enable()
Sets all predicates to true.
Definition: predicate_vector.h:435
cutlass::PredicateVector::ConstIterator::operator==
CUTLASS_HOST_DEVICE bool operator==(ConstIterator const &it) const
Returns true if iterators point to the same bit.
Definition: predicate_vector.h:362
cutlass::PredicateVector::operator[]
CUTLASS_HOST_DEVICE bool operator[](int idx) const
Accesses a bit within the predicate vector.
Definition: predicate_vector.h:443
cutlass::PredicateVector::Iterator::operator==
CUTLASS_HOST_DEVICE bool operator==(Iterator const &it) const
Returns true if iterators point to the same bit.
Definition: predicate_vector.h:256
cutlass::PredicateVector::Iterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(Iterator const &it) const
Returns false if iterators point to the same bit.
Definition: predicate_vector.h:260
cutlass::PredicateVector::at
CUTLASS_HOST_DEVICE bool at(int idx) const
Accesses a bit within the predicate vector.
Definition: predicate_vector.h:446
cutlass::PredicateVector::Iterator::Iterator
CUTLASS_HOST_DEVICE Iterator(PredicateVector &vec, int _start=0)
Constructs an iterator from a PredicateVector.
Definition: predicate_vector.h:192
cutlass::PredicateVector::ConstIterator::operator++
CUTLASS_HOST_DEVICE ConstIterator operator++(int)
Post-increment.
Definition: predicate_vector.h:330
cutlass::PredicateVector::Iterator::operator++
CUTLASS_HOST_DEVICE Iterator operator++(int)
Post-increment.
Definition: predicate_vector.h:224
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator(Iterator const &it)
Copy constructor.
Definition: predicate_vector.h:389
cutlass::PredicateVector::Iterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:272
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
cutlass::PredicateVector::TrivialIterator
Iterator that always returns true.
Definition: predicate_vector.h:382
cutlass::PredicateVector::TrivialIterator::operator++
CUTLASS_HOST_DEVICE TrivialIterator operator++(int)
Post-increment.
Definition: predicate_vector.h:401
cutlass::PredicateVector::ConstIterator::operator+
CUTLASS_HOST_DEVICE ConstIterator operator+(int offset)
Iterator advances by some amount.
Definition: predicate_vector.h:346
cutlass::PredicateVector::ConstIterator::operator+=
CUTLASS_HOST_DEVICE ConstIterator & operator+=(int offset)
Increment.
Definition: predicate_vector.h:309
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: cutlass.h:110
cutlass::PredicateVector::const_begin
CUTLASS_DEVICE ConstIterator const_begin() const
Returns a ConstIterator.
Definition: predicate_vector.h:506
cutlass::PredicateVector::Iterator::operator--
CUTLASS_HOST_DEVICE Iterator & operator--()
Pre-decrement.
Definition: predicate_vector.h:210
cutlass::PredicateVector::Iterator::operator-=
CUTLASS_HOST_DEVICE Iterator & operator-=(int offset)
Decrement.
Definition: predicate_vector.h:217
cutlass::PredicateVector::Iterator::Iterator
CUTLASS_HOST_DEVICE Iterator(Iterator const &it)
Copy constructor.
Definition: predicate_vector.h:188
cutlass::PredicateVector::ConstIterator::operator-
CUTLASS_HOST_DEVICE ConstIterator operator-(int offset)
Iterator recedes by some amount.
Definition: predicate_vector.h:354
cutlass::PredicateVector::Iterator::operator+
CUTLASS_HOST_DEVICE Iterator operator+(int offset)
Iterator advances by some amount.
Definition: predicate_vector.h:240
cutlass::PredicateVector::fill
CUTLASS_HOST_DEVICE void fill(bool value=true)
Fills all predicates with a given value.
Definition: predicate_vector.h:417
cutlass::PredicateVector::kPredicates
static int const kPredicates
Number of bits stored by the PredicateVector.
Definition: predicate_vector.h:108
cutlass::PredicateVector::end
CUTLASS_DEVICE Iterator end()
Returns an iterator.
Definition: predicate_vector.h:502
CUTLASS_ASSERT
#define CUTLASS_ASSERT(x)
Definition: cutlass.h:92
cutlass::PredicateVector::Iterator::operator+=
CUTLASS_HOST_DEVICE Iterator & operator+=(int offset)
Increment.
Definition: predicate_vector.h:203
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::PredicateVector::kPredicatesPerByte
static int const kPredicatesPerByte
Number of bits stored within each byte of the predicate bit vector.
Definition: predicate_vector.h:111
cutlass::PredicateVector::operator&=
CUTLASS_HOST_DEVICE PredicateVector & operator&=(PredicateVector const &predicates)
Computes the intersection of two identical predicate vectors.
Definition: predicate_vector.h:465
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::PredicateVector
Statically sized array of bits implementing.
Definition: predicate_vector.h:106
cutlass::PredicateVector::kWordCount
static int const kWordCount
Number of storage elements needed.
Definition: predicate_vector.h:129
cutlass::PredicateVector::ConstIterator::operator!=
CUTLASS_HOST_DEVICE bool operator!=(ConstIterator const &it) const
Returns false if iterators point to the same bit.
Definition: predicate_vector.h:366
cutlass::PredicateVector::Iterator::at
CUTLASS_HOST_DEVICE bool at() const
Gets the bit at the pointed to location.
Definition: predicate_vector.h:268
cutlass::PredicateVector::Iterator::operator++
CUTLASS_HOST_DEVICE Iterator & operator++()
Pre-increment.
Definition: predicate_vector.h:196
cutlass::PredicateVector::ConstIterator
An iterator implementing Predicate Iterator Concept enabling sequential read and write access to pred...
Definition: predicate_vector.h:284
cutlass::PredicateVector::set
CUTLASS_HOST_DEVICE void set(int idx, bool value=true)
Set a bit within the predicate vector.
Definition: predicate_vector.h:454
cutlass::PredicateVector::Iterator::operator-
CUTLASS_HOST_DEVICE Iterator operator-(int offset)
Iterator recedes by some amount.
Definition: predicate_vector.h:248
cutlass::PredicateVector::kPredicateStart
static int const kPredicateStart
First bit withing each byte containing predicates.
Definition: predicate_vector.h:114
cutlass::PredicateVector::ConstIterator::ConstIterator
CUTLASS_HOST_DEVICE ConstIterator(ConstIterator const &it)
Copy constructor.
Definition: predicate_vector.h:294
cutlass::PredicateVector::ConstIterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:378
cutlass::PredicateVector::ConstIterator::operator--
CUTLASS_HOST_DEVICE ConstIterator operator--(int)
Post-decrement.
Definition: predicate_vector.h:338
cutlass::PredicateVector::clear
CUTLASS_HOST_DEVICE void clear()
Clears all predicates.
Definition: predicate_vector.h:427
cutlass::PredicateVector::PredicateVector
CUTLASS_HOST_DEVICE PredicateVector(bool value=true)
Initialize the predicate vector.
Definition: predicate_vector.h:414
cutlass::PredicateVector::begin
CUTLASS_DEVICE Iterator begin()
Returns an iterator to the start of the bit vector.
Definition: predicate_vector.h:498
cutlass.h
Basic include for CUTLASS.
cutlass::PredicateVector::TrivialIterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:405
cutlass::PredicateVector::Iterator
An iterator implementing Predicate Iterator Concept enabling sequential read and write access to pred...
Definition: predicate_vector.h:178
cutlass::PredicateVector::Iterator::operator--
CUTLASS_HOST_DEVICE Iterator operator--(int)
Post-decrement.
Definition: predicate_vector.h:232
cutlass::PredicateVector::ConstIterator::operator-=
CUTLASS_HOST_DEVICE ConstIterator & operator-=(int offset)
Decrement.
Definition: predicate_vector.h:323
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator()
Constructor.
Definition: predicate_vector.h:385