Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

predicate_vector.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2018, 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 **************************************************************************************************/
29 #pragma once
30
31 #include <assert.h>
32 #include <stdint.h>
33
34 #include "cutlass/cutlass.h"
35 #include "cutlass/shape.h"
36
37 #include "cutlass/util/platform.h"
38
39 namespace cutlass {
40
42
59
79
95
98 template <
100 int kPredicates_,
102 int kPredicatesPerByte_ = 4,
104 int kPredicateStart_ = 0>
105 struct PredicateVector {
107 static int const kPredicates = kPredicates_;
108
110 static int const kPredicatesPerByte = kPredicatesPerByte_;
111
113 static int const kPredicateStart = kPredicateStart_;
114
115 // Make sure no one tries to put more than 8 bits in a byte :)
116 static_assert(kPredicatesPerByte <= 8, "kPredicatesPerByte must fit within an actual byte");
117 // Make sure the "offsetted" bits fit in one byte.
118 static_assert(kPredicateStart + kPredicatesPerByte <= 8,
119 "The offsetted predicates must fit within an actual byte.");
120
122 typedef uint32_t Storage;
123
125 static int const kBytes = (kPredicates + kPredicatesPerByte - 1) / kPredicatesPerByte;
126
128 static int const kWordCount = (kBytes + sizeof(Storage) - 1) / sizeof(Storage);
129
130 private:
131 //
132 // Data members
133 //
134
136 Storage storageData[kWordCount];
137
138 //
139 // Methods
140 //
141
143 CUTLASS_HOST_DEVICE void computeStorageOffset(int &word, int &bit, int idx) const {
144 CUTLASS_ASSERT(idx < kPredicates);
145
146 int byte = (idx / kPredicatesPerByte);
147 int bit_offset = (idx % kPredicatesPerByte);
148
149 word = byte / sizeof(Storage);
150 int byte_offset = (byte % sizeof(Storage));
151
152 bit = byte_offset * 8 + bit_offset + kPredicateStart;
153 }
154
156 CUTLASS_HOST_DEVICE Storage &storage(int word) {
157 CUTLASS_ASSERT(word < kWordCount);
158 return storageData[word];
159 }
160
162 CUTLASS_HOST_DEVICE Storage const &storage(int word) const {
163 CUTLASS_ASSERT(word < kWordCount);
164 return storageData[word];
165 }
166
167 public:
168 //
169 // Iterator
170 //
171
177 class ConstIterator {
179 PredicateVector const &vec_;
180
182 int bit_;
183
184 public:
186 CUTLASS_HOST_DEVICE
187 ConstIterator(ConstIterator const &it) : vec_(it.vec_), bit_(it.bit_) {}
188
190 CUTLASS_HOST_DEVICE
191 ConstIterator(PredicateVector const &_vec, int _start = 0) : vec_(_vec), bit_(_start) {}
192
194 CUTLASS_HOST_DEVICE
195 ConstIterator &operator++() {
196 ++bit_;
197 return *this;
198 }
199
201 CUTLASS_HOST_DEVICE
202 ConstIterator &operator--() {
203 --bit_;
204 return *this;
205 }
206
208 CUTLASS_HOST_DEVICE
209 ConstIterator operator++(int) {
210 ConstIterator ret(*this);
211 ret.bit_++;
212 return ret;
213 }
214
216 CUTLASS_HOST_DEVICE
217 ConstIterator operator--(int) {
218 ConstIterator ret(*this);
219 ret.bit_--;
220 return ret;
221 }
222
224 CUTLASS_HOST_DEVICE
225 bool operator==(ConstIterator const &it) const { return bit_ == it.bit_; }
226
228 CUTLASS_HOST_DEVICE
229 bool operator!=(ConstIterator const &it) const { return bit_ != it.bit_; }
230
232 CUTLASS_HOST_DEVICE
233 bool operator*() const { return vec_[bit_]; }
234 };
235
241 class Iterator {
243 PredicateVector &vec_;
244
246 int bit_;
247
248 public:
250 CUTLASS_HOST_DEVICE
251 Iterator(Iterator const &it) : vec_(it.vec_), bit_(it.bit_) {}
252
254 CUTLASS_HOST_DEVICE
255 Iterator(PredicateVector &_vec, int _start = 0) : vec_(_vec), bit_(_start) {}
256
258 CUTLASS_HOST_DEVICE
259 Iterator &operator++() {
260 ++bit_;
261 return *this;
262 }
263
265 CUTLASS_HOST_DEVICE
266 Iterator &operator--() {
267 --bit_;
268 return *this;
269 }
270
272 CUTLASS_HOST_DEVICE
273 Iterator operator++(int) {
274 Iterator ret(*this);
275 ret.bit_++;
276 return ret;
277 }
278
280 CUTLASS_HOST_DEVICE
281 Iterator operator--(int) {
282 Iterator ret(*this);
283 ret.bit_--;
284 return ret;
285 }
286
288 CUTLASS_HOST_DEVICE
289 bool operator==(Iterator const &it) const { return bit_ == it.bit_; }
290
292 CUTLASS_HOST_DEVICE
293 bool operator!=(Iterator const &it) const { return bit_ != it.bit_; }
294
296 CUTLASS_HOST_DEVICE
297 bool get() { return vec_[bit_]; }
298
300 CUTLASS_HOST_DEVICE
301 bool operator*() const { return vec_[bit_]; }
302
304 CUTLASS_HOST_DEVICE
305 void set(bool value = true) { vec_.set(bit_, value); }
306 };
307
309 struct TrivialIterator {
311 CUTLASS_HOST_DEVICE
312 TrivialIterator() {}
313
315 CUTLASS_HOST_DEVICE
316 TrivialIterator(Iterator const &it) {}
317
319 CUTLASS_HOST_DEVICE
320 TrivialIterator(PredicateVector const &_vec) {}
321
323 CUTLASS_HOST_DEVICE
324 TrivialIterator &operator++() { return *this; }
325
327 CUTLASS_HOST_DEVICE
328 TrivialIterator operator++(int) { return *this; }
329
331 CUTLASS_HOST_DEVICE
332 bool operator*() const { return true; }
333 };
334
335 public:
336 //
337 // Methods
338 //
339
341 CUTLASS_HOST_DEVICE PredicateVector(bool value = true) { fill(value); }
342
344 CUTLASS_HOST_DEVICE void fill(bool value = true) {
345 Storage item = (value ? ~Storage(0) : Storage(0));
346
347 CUTLASS_PRAGMA_UNROLL
348 for (int i = 0; i < kWordCount; ++i) {
349 storage(i) = item;
350 }
351 }
352
354 CUTLASS_HOST_DEVICE bool operator[](int idx) const { return at(idx); }
355
357 CUTLASS_HOST_DEVICE bool at(int idx) const {
358 int bit, word;
359 computeStorageOffset(word, bit, idx);
360
361 return ((storage(word) >> bit) & 1);
362 }
363
365 CUTLASS_HOST_DEVICE void set(int idx, bool value = true) {
366 int bit, word;
367 computeStorageOffset(word, bit, idx);
368
369 Storage disable_mask = (~(Storage(1) << bit));
370 Storage enable_mask = (Storage(value) << bit);
371
372 storage(word) = ((storage(word) & disable_mask) | enable_mask);
373 }
374
376 CUTLASS_HOST_DEVICE PredicateVector &operator&=(PredicateVector const &predicates) {
377 CUTLASS_PRAGMA_UNROLL
378 for (int i = 0; i < kWordCount; ++i) {
379 storage(i) = (storage(i) & predicates.storage(i));
380 }
381 return *this;
382 }
383
385 CUTLASS_HOST_DEVICE PredicateVector &operator|=(PredicateVector const &predicates) {
386 CUTLASS_PRAGMA_UNROLL
387 for (int i = 0; i < kWordCount; ++i) {
388 storage(i) = (storage(i) | predicates.storage(i));
389 }
390 return *this;
391 }
392
394 CUTLASS_HOST_DEVICE bool is_zero() const {
395 Storage mask(0);
396 for (int byte = 0; byte < sizeof(Storage); ++byte) {
397 Storage byte_mask = (((1 << kPredicatesPerByte) - 1) << kPredicateStart);
398 mask |= (byte_mask << (byte * 8));
399 }
400 uint32_t result = 0;
401 for (int word = 0; word < kWordCount; ++word) {
402 result |= storage(word);
403 }
404 return result == 0;
405 }
406
408 CUTLASS_DEVICE
409 Iterator begin() { return Iterator(*this); }
410
412 CUTLASS_DEVICE
413 Iterator end() { return Iterator(*this, kPredicates); }
414
416 CUTLASS_DEVICE
417 ConstIterator const_begin() const { return ConstIterator(*this); }
418
420 CUTLASS_DEVICE
421 ConstIterator const_end() const { return ConstIterator(*this, kPredicates); }
422 };
423
425
427 struct TrivialPredicateTileAdapter {
429 CUTLASS_HOST_DEVICE TrivialPredicateTileAdapter() {}
430
432 CUTLASS_HOST_DEVICE bool at(int, int, int, int) const { return true; }
433 };
434
436
438 template <typename PredicateVector_, typename Iterations_>
439 struct PredicateTileAdapter {
441 typedef PredicateVector_ PredicateVector;
443 typedef Iterations_ Iterations;
444
445 private:
447 PredicateVector &predicates;
448
449 public:
451 CUTLASS_DEVICE PredicateTileAdapter(PredicateVector &predicates_) : predicates(predicates_) {}
452
454 CUTLASS_DEVICE bool at(int d, int h, int w, int c) const {
455 int const bit = ComputeOffsetFromShape<Iterations>::get(d, h, w, c);
456 return predicates.at(bit);
457 }
458
460 CUTLASS_DEVICE void set(int d, int h, int w, int c, bool value) {
461 int const bit = ComputeOffsetFromShape<Iterations>::get(d, h, w, c);
462 predicates.set(bit, value);
463 }
464 };
465
467
469 template <typename PredicateVector_, typename Iterations_>
470 struct ConstPredicateTileAdapter {
472 typedef PredicateVector_ PredicateVector;
474 typedef Iterations_ Iterations;
475
476 private:
478 PredicateVector const &predicates;
479
480 public:
482 CUTLASS_DEVICE ConstPredicateTileAdapter(PredicateVector const &predicates_)
483 : predicates(predicates_) {}
484
486 CUTLASS_DEVICE bool at(int d, int h, int w, int c) const {
487 int const bit = ComputeOffsetFromShape<Iterations>::get(d, h, w, c);
488 return predicates.at(bit);
489 }
490 };
491
493
494 } // namespace cutlass
cutlass::PredicateVector::Iterator::Iterator
CUTLASS_HOST_DEVICE Iterator(PredicateVector &_vec, int _start=0)
Constructs an iterator from a PredicateVector.
Definition: predicate_vector.h:255
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:229
cutlass::PredicateVector::operator|=
CUTLASS_HOST_DEVICE PredicateVector & operator|=(PredicateVector const &predicates)
Computes the union of two identical predicate vectors.
Definition: predicate_vector.h:385
cutlass::PredicateVector::TrivialIterator::operator++
CUTLASS_HOST_DEVICE TrivialIterator & operator++()
Pre-increment.
Definition: predicate_vector.h:324
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: performance_tuning.h:35
cutlass
Definition: convert.h:33
cutlass::PredicateVector::is_zero
CUTLASS_HOST_DEVICE bool is_zero() const
Returns true if entire predicate array is zero.
Definition: predicate_vector.h:394
cutlass::PredicateVector::Storage
uint32_t Storage
Storage type of individual elements.
Definition: predicate_vector.h:116
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator(PredicateVector const &_vec)
Constructs an iterator from a PredicateVector.
Definition: predicate_vector.h:320
cutlass::PredicateVector::ConstIterator::operator--
CUTLASS_HOST_DEVICE ConstIterator & operator--()
Pre-decrement.
Definition: predicate_vector.h:202
cutlass::PredicateVector::kBytes
static int const kBytes
Number of bytes needed.
Definition: predicate_vector.h:125
cutlass::PredicateVector::const_begin
CUTLASS_DEVICE ConstIterator const_begin() const
Returns a ConstIterator.
Definition: predicate_vector.h:417
cutlass::PredicateVector::ConstIterator::ConstIterator
CUTLASS_HOST_DEVICE ConstIterator(PredicateVector const &_vec, int _start=0)
Definition: predicate_vector.h:191
cutlass::PredicateVector::at
CUTLASS_HOST_DEVICE bool at(int idx) const
Accesses a bit within the predicate vector.
Definition: predicate_vector.h:357
cutlass::PredicateVector::ConstIterator::operator++
CUTLASS_HOST_DEVICE ConstIterator & operator++()
Pre-increment.
Definition: predicate_vector.h:195
cutlass::PredicateTileAdapter::PredicateVector
PredicateVector_ PredicateVector
The vector of predicates.
Definition: predicate_vector.h:441
cutlass::PredicateVector::ConstIterator::operator++
CUTLASS_HOST_DEVICE ConstIterator operator++(int)
Post-increment.
Definition: predicate_vector.h:209
cutlass::PredicateVector::Iterator::operator++
CUTLASS_HOST_DEVICE Iterator operator++(int)
Post-increment.
Definition: predicate_vector.h:273
cutlass::PredicateTileAdapter
Adapter to enable random access to predicates via logical coordinate within a tile.
Definition: predicate_vector.h:439
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator(Iterator const &it)
Copy constructor.
Definition: predicate_vector.h:316
cutlass::ComputeOffsetFromShape::get
static CUTLASS_HOST_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:181
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:309
cutlass::PredicateVector::TrivialIterator::operator++
CUTLASS_HOST_DEVICE TrivialIterator operator++(int)
Post-increment.
Definition: predicate_vector.h:328
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:289
cutlass::PredicateTileAdapter::PredicateTileAdapter
CUTLASS_DEVICE PredicateTileAdapter(PredicateVector &predicates_)
Ctor.
Definition: predicate_vector.h:451
cutlass::PredicateTileAdapter::at
CUTLASS_DEVICE bool at(int d, int h, int w, int c) const
Get the value at location (d, h, w, c).
Definition: predicate_vector.h:454
cutlass::ConstPredicateTileAdapter::at
CUTLASS_DEVICE bool at(int d, int h, int w, int c) const
Get the value at location (d, h, w, c).
Definition: predicate_vector.h:486
cutlass::PredicateVector::Iterator::operator--
CUTLASS_HOST_DEVICE Iterator & operator--()
Pre-decrement.
Definition: predicate_vector.h:266
cutlass::ConstPredicateTileAdapter::PredicateVector
PredicateVector_ PredicateVector
The vector of predicates.
Definition: predicate_vector.h:472
cutlass::PredicateVector::operator &=
CUTLASS_HOST_DEVICE PredicateVector & operator &=(PredicateVector const &predicates)
Computes the intersection of two identical predicate vectors.
Definition: predicate_vector.h:376
cutlass::PredicateVector::Iterator::Iterator
CUTLASS_HOST_DEVICE Iterator(Iterator const &it)
Copy constructor.
Definition: predicate_vector.h:251
cutlass::PredicateVector::operator[]
CUTLASS_HOST_DEVICE bool operator[](int idx) const
Accesses a bit within the predicate vector.
Definition: predicate_vector.h:354
cutlass::PredicateVector::Iterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:301
cutlass::PredicateVector::TrivialIterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:332
cutlass::PredicateVector::fill
CUTLASS_HOST_DEVICE void fill(bool value=true)
Fills all predicates with a given value.
Definition: predicate_vector.h:344
cutlass::PredicateVector::kPredicates
static int const kPredicates
Number of bits stored by the PredicateVector.
Definition: predicate_vector.h:107
cutlass::PredicateVector::end
CUTLASS_DEVICE Iterator end()
Returns an iterator.
Definition: predicate_vector.h:413
CUTLASS_ASSERT
#define CUTLASS_ASSERT(x)
Definition: cutlass.h:50
cutlass::TrivialPredicateTileAdapter::at
CUTLASS_HOST_DEVICE bool at(int, int, int, int) const
The value at location (d, h, w, c).
Definition: predicate_vector.h:432
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::PredicateVector::kPredicatesPerByte
static int const kPredicatesPerByte
Number of bits stored within each byte of the predicate bit vector.
Definition: predicate_vector.h:110
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::PredicateVector
Statically sized array of bits implementing.
Definition: predicate_vector.h:105
cutlass::PredicateVector::kWordCount
static int const kWordCount
Number of storage elements needed.
Definition: predicate_vector.h:128
cutlass::PredicateVector::const_end
CUTLASS_DEVICE ConstIterator const_end() const
Returns a ConstIterator.
Definition: predicate_vector.h:421
cutlass::TrivialPredicateTileAdapter
Always returns true predicate.
Definition: predicate_vector.h:427
cutlass::PredicateVector::Iterator::operator++
CUTLASS_HOST_DEVICE Iterator & operator++()
Pre-increment.
Definition: predicate_vector.h:259
cutlass::PredicateVector::ConstIterator
A const iterator implementing Predicate Iterator Concept enabling sequential read-only access to pred...
Definition: predicate_vector.h:177
cutlass::PredicateVector::set
CUTLASS_HOST_DEVICE void set(int idx, bool value=true)
Set a bit within the predicate vector.
Definition: predicate_vector.h:365
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:225
cutlass::ConstPredicateTileAdapter::Iterations
Iterations_ Iterations
The iterations.
Definition: predicate_vector.h:474
cutlass::PredicateTileAdapter::Iterations
Iterations_ Iterations
The iterations.
Definition: predicate_vector.h:443
cutlass::PredicateVector::ConstIterator::operator*
CUTLASS_HOST_DEVICE bool operator*() const
Dereferences iterator.
Definition: predicate_vector.h:233
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:293
cutlass::PredicateVector::kPredicateStart
static int const kPredicateStart
First bit withing each byte containing predicates.
Definition: predicate_vector.h:113
cutlass::PredicateVector::ConstIterator::ConstIterator
CUTLASS_HOST_DEVICE ConstIterator(ConstIterator const &it)
Copy constructor.
Definition: predicate_vector.h:187
cutlass::TrivialPredicateTileAdapter::TrivialPredicateTileAdapter
CUTLASS_HOST_DEVICE TrivialPredicateTileAdapter()
Ctor.
Definition: predicate_vector.h:429
cutlass::PredicateVector::ConstIterator::operator--
CUTLASS_HOST_DEVICE ConstIterator operator--(int)
Post-decrement.
Definition: predicate_vector.h:217
cutlass::ConstPredicateTileAdapter
Adapter to enable random access to predicates via logical coordinate within a tile.
Definition: predicate_vector.h:470
cutlass::ConstPredicateTileAdapter::ConstPredicateTileAdapter
CUTLASS_DEVICE ConstPredicateTileAdapter(PredicateVector const &predicates_)
Ctor.
Definition: predicate_vector.h:482
shape.h
Defines Shape implementing the Layout concept for representing a 4D hypercube of objects.
cutlass::PredicateVector::PredicateVector
CUTLASS_HOST_DEVICE PredicateVector(bool value=true)
Initialize the predicate vector.
Definition: predicate_vector.h:341
cutlass::PredicateVector::begin
CUTLASS_DEVICE Iterator begin()
Returns an iterator to the start of the bit vector.
Definition: predicate_vector.h:409
cutlass.h
Basic include for CUTLASS macros.
cutlass::PredicateVector::Iterator
An iterator implementing Predicate Iterator Concept enabling sequential read and write access to pred...
Definition: predicate_vector.h:241
cutlass::PredicateVector::Iterator::operator--
CUTLASS_HOST_DEVICE Iterator operator--(int)
Post-decrement.
Definition: predicate_vector.h:281
cutlass::PredicateVector::TrivialIterator::TrivialIterator
CUTLASS_HOST_DEVICE TrivialIterator()
Constructor.
Definition: predicate_vector.h:312