| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

array_subbyte.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2019, 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 **************************************************************************************************/
30 #pragma once
31
32 #include "cutlass/cutlass.h"
33 #include "cutlass/array.h"
34 #include "cutlass/platform/platform.h"
35
36 namespace cutlass {
37
39
41 template <
42 typename T,
43 int N
44 >
45 class Array<T, N, false> {
46 public:
47
48 static_assert(sizeof_bits<T>::value * N >= 8,
49 "Array<> specialized for sub-byte types assume the actual stored element size is 1 byte");
50
51 static int const kSizeBits = sizeof_bits<T>::value * N;
52
54 using Storage = typename platform::conditional<
55 ((kSizeBits % 32) != 0),
56 typename platform::conditional<
57 ((kSizeBits % 16) != 0),
58 uint8_t,
59 uint16_t
60 >::type,
61 uint32_t
62 >::type;
63
65 using Element = T;
66
68 static int const kElementsPerStoredItem = (sizeof(Storage) * 8) / sizeof_bits<T>::value;
69
71 static size_t const kStorageElements = N / kElementsPerStoredItem;
72
74 static size_t const kElements = N;
75
77 static Storage const kMask = ((Storage(1) << sizeof_bits<T>::value) - 1);
78
79 //
80 // C++ standard members with pointer types removed
81 //
82
83 typedef T value_type;
84 typedef size_t size_type;
85 typedef ptrdiff_t difference_type;
86 typedef value_type *pointer;
87 typedef value_type const *const_pointer;
88
89 //
90 // References
91 //
92
94 class reference {
96 Storage *ptr_;
97
99 int idx_;
100
101 public:
102
104 CUTLASS_HOST_DEVICE
105 reference(): ptr_(nullptr), idx_(0) { }
106
108 CUTLASS_HOST_DEVICE
109 reference(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
110
112 CUTLASS_HOST_DEVICE
113 reference &operator=(T x) {
114 Storage item = (reinterpret_cast<Storage const &>(x) & kMask);
115
116 Storage kUpdateMask = Storage(~(kMask << (idx_ * sizeof_bits<T>::value)));
117 *ptr_ = Storage(((*ptr_ & kUpdateMask) | (item << idx_ * sizeof_bits<T>::value)));
118
119 return *this;
120 }
121
122 CUTLASS_HOST_DEVICE
123 T get() const {
124 Storage item = Storage((*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask);
125 return reinterpret_cast<T const &>(item);
126 }
127
129 CUTLASS_HOST_DEVICE
130 operator T() const {
131 return get();
132 }
133
135 CUTLASS_HOST_DEVICE
136 explicit operator int() const {
137 return int(get());
138 }
139
141 CUTLASS_HOST_DEVICE
142 explicit operator float() const {
143 return float(get());
144 }
145 };
146
148 class const_reference {
149
151 Storage const *ptr_;
152
154 int idx_;
155
156 public:
157
159 CUTLASS_HOST_DEVICE
160 const_reference(): ptr_(nullptr), idx_(0) { }
161
163 CUTLASS_HOST_DEVICE
164 const_reference(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
165
166 CUTLASS_HOST_DEVICE
167 const T get() const {
168 Storage item = (*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask;
169 return reinterpret_cast<T const &>(item);
170 }
171
173 CUTLASS_HOST_DEVICE
174 operator T() const {
175 Storage item = Storage(Storage(*ptr_ >> Storage(idx_ * sizeof_bits<T>::value)) & kMask);
176 return reinterpret_cast<T const &>(item);
177 }
178
180 CUTLASS_HOST_DEVICE
181 explicit operator int() const {
182 return int(get());
183 }
184
186 CUTLASS_HOST_DEVICE
187 explicit operator float() const {
188 return float(get());
189 }
190 };
191
192 //
193 // Iterators
194 //
195
197 class iterator {
198
200 Storage *ptr_;
201
203 int idx_;
204
205 public:
206
207 CUTLASS_HOST_DEVICE
208 iterator(): ptr_(nullptr), idx_(0) { }
209
210 CUTLASS_HOST_DEVICE
211 iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
212
213 CUTLASS_HOST_DEVICE
214 iterator &operator++() {
215 ++idx_;
216 if (idx_ == kElementsPerStoredItem) {
217 ++ptr_;
218 idx_ = 0;
219 }
220 return *this;
221 }
222
223 CUTLASS_HOST_DEVICE
224 iterator &operator--() {
225 if (!idx_) {
226 --ptr_;
227 idx_ = kElementsPerStoredItem - 1;
228 }
229 else {
230 --idx_;
231 }
232 return *this;
233 }
234
235 CUTLASS_HOST_DEVICE
236 iterator operator++(int) {
237 iterator ret(*this);
238 ++idx_;
239 if (idx_ == kElementsPerStoredItem) {
240 ++ptr_;
241 idx_ = 0;
242 }
243 return ret;
244 }
245
246 CUTLASS_HOST_DEVICE
247 iterator operator--(int) {
248 iterator ret(*this);
249 if (!idx_) {
250 --ptr_;
251 idx_ = kElementsPerStoredItem - 1;
252 }
253 else {
254 --idx_;
255 }
256 return ret;
257 }
258
259 CUTLASS_HOST_DEVICE
260 reference operator*() const {
261 return reference(ptr_, idx_);
262 }
263
264 CUTLASS_HOST_DEVICE
265 bool operator==(iterator const &other) const {
266 return ptr_ == other.ptr_ && idx_ == other.idx_;
267 }
268
269 CUTLASS_HOST_DEVICE
270 bool operator!=(iterator const &other) const {
271 return !(*this == other);
272 }
273 };
274
276 class const_iterator {
277
279 Storage const *ptr_;
280
282 int idx_;
283
284 public:
285
286 CUTLASS_HOST_DEVICE
287 const_iterator(): ptr_(nullptr), idx_(0) { }
288
289 CUTLASS_HOST_DEVICE
290 const_iterator(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
291
292 CUTLASS_HOST_DEVICE
293 iterator &operator++() {
294 ++idx_;
295 if (idx_ == kElementsPerStoredItem) {
296 ++ptr_;
297 idx_ = 0;
298 }
299 return *this;
300 }
301
302 CUTLASS_HOST_DEVICE
303 iterator &operator--() {
304 if (!idx_) {
305 --ptr_;
306 idx_ = kElementsPerStoredItem - 1;
307 }
308 else {
309 --idx_;
310 }
311 return *this;
312 }
313
314 CUTLASS_HOST_DEVICE
315 iterator operator++(int) {
316 iterator ret(*this);
317 ++idx_;
318 if (idx_ == kElementsPerStoredItem) {
319 ++ptr_;
320 idx_ = 0;
321 }
322 return ret;
323 }
324
325 CUTLASS_HOST_DEVICE
326 iterator operator--(int) {
327 iterator ret(*this);
328 if (!idx_) {
329 --ptr_;
330 idx_ = kElementsPerStoredItem - 1;
331 }
332 else {
333 --idx_;
334 }
335 return ret;
336 }
337
338 CUTLASS_HOST_DEVICE
339 const_reference operator*() const {
340 return const_reference(ptr_, idx_);
341 }
342
343 CUTLASS_HOST_DEVICE
344 bool operator==(iterator const &other) const {
345 return ptr_ == other.ptr_ && idx_ == other.idx_;
346 }
347
348 CUTLASS_HOST_DEVICE
349 bool operator!=(iterator const &other) const {
350 return !(*this == other);
351 }
352 };
353
355 class reverse_iterator {
356
358 Storage *ptr_;
359
361 int idx_;
362
363 public:
364
365 CUTLASS_HOST_DEVICE
366 reverse_iterator(): ptr_(nullptr), idx_(0) { }
367
368 CUTLASS_HOST_DEVICE
369 reverse_iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
370
371 // TODO
372 };
373
375 class const_reverse_iterator {
376
378 Storage const *ptr_;
379
381 int idx_;
382
383 public:
384
385 CUTLASS_HOST_DEVICE
386 const_reverse_iterator(): ptr_(nullptr), idx_(0) { }
387
388 CUTLASS_HOST_DEVICE
389 const_reverse_iterator(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }
390
391 // TODO
392 };
393
394 private:
395
397 Storage storage[kStorageElements];
398
399 public:
400
401 CUTLASS_HOST_DEVICE
402 Array() { }
403
404 CUTLASS_HOST_DEVICE
405 Array(Array const &x) {
406 CUTLASS_PRAGMA_UNROLL
407 for (int i = 0; i < int(kStorageElements); ++i) {
408 storage[i] = x.storage[i];
409 }
410 }
411
413 CUTLASS_HOST_DEVICE
414 void clear() {
415
416 CUTLASS_PRAGMA_UNROLL
417 for (int i = 0; i < int(kStorageElements); ++i) {
418 storage[i] = Storage(0);
419 }
420 }
421
422 CUTLASS_HOST_DEVICE
423 reference at(size_type pos) {
424 return reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);
425 }
426
427 CUTLASS_HOST_DEVICE
428 const_reference at(size_type pos) const {
429 return const_reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);
430 }
431
432 CUTLASS_HOST_DEVICE
433 reference operator[](size_type pos) {
434 return at(pos);
435 }
436
437 CUTLASS_HOST_DEVICE
438 const_reference operator[](size_type pos) const {
439 return at(pos);
440 }
441
442 CUTLASS_HOST_DEVICE
443 reference front() {
444 return at(0);
445 }
446
447 CUTLASS_HOST_DEVICE
448 const_reference front() const {
449 return at(0);
450 }
451
452 CUTLASS_HOST_DEVICE
453 reference back() {
454 return reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);
455 }
456
457 CUTLASS_HOST_DEVICE
458 const_reference back() const {
459 return const_reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);
460 }
461
462 CUTLASS_HOST_DEVICE
463 pointer data() {
464 return reinterpret_cast<pointer>(storage);
465 }
466
467 CUTLASS_HOST_DEVICE
468 const_pointer data() const {
469 return reinterpret_cast<const_pointer>(storage);
470 }
471
472 CUTLASS_HOST_DEVICE
473 Storage * raw_data() {
474 return storage;
475 }
476
477 CUTLASS_HOST_DEVICE
478 Storage const * raw_data() const {
479 return storage;
480 }
481
482
483 CUTLASS_HOST_DEVICE
484 constexpr bool empty() const {
485 return !kElements;
486 }
487
488 CUTLASS_HOST_DEVICE
489 constexpr size_type size() const {
490 return kElements;
491 }
492
493 CUTLASS_HOST_DEVICE
494 constexpr size_type max_size() const {
495 return kElements;
496 }
497
498 CUTLASS_HOST_DEVICE
499 void fill(T const &value) {
500 // TODO
501 }
502
503 CUTLASS_HOST_DEVICE
504 iterator begin() {
505 return iterator(storage);
506 }
507
508 CUTLASS_HOST_DEVICE
509 const_iterator cbegin() const {
510 return const_iterator(storage);
511 }
512
513 CUTLASS_HOST_DEVICE
514 iterator end() {
515 return iterator(storage + kStorageElements);
516 }
517
518 CUTLASS_HOST_DEVICE
519 const_iterator cend() const {
520 return const_iterator(storage + kStorageElements);
521 }
522
523 CUTLASS_HOST_DEVICE
524 reverse_iterator rbegin() {
525 return reverse_iterator(storage + kStorageElements);
526 }
527
528 CUTLASS_HOST_DEVICE
529 const_reverse_iterator crbegin() const {
530 return const_reverse_iterator(storage + kStorageElements);
531 }
532
533 CUTLASS_HOST_DEVICE
534 reverse_iterator rend() {
535 return reverse_iterator(storage);
536 }
537
538 CUTLASS_HOST_DEVICE
539 const_reverse_iterator crend() const {
540 return const_reverse_iterator(storage);
541 }
542
543 //
544 // Comparison operators
545 //
546
547 };
548
550
551 } // namespace cutlass
552
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