| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

regular_tile_access_iterator_tensor_op.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/array.h"
33 #include "cutlass/cutlass.h"
34 #include "cutlass/layout/pitch_linear.h"
35 #include "cutlass/layout/tensor_op_multiplicand_sm75.h"
36 #include "cutlass/matrix_coord.h"
37 #include "cutlass/matrix_shape.h"
38 #include "cutlass/tensor_ref.h"
39 #include "cutlass/transform/threadblock/regular_tile_access_iterator.h"
40
42
43 namespace cutlass {
44 namespace transform {
45 namespace threadblock {
46
48
56 template <typename Shape_, typename Element_, int AdvanceRank,
57 typename ThreadMap_, int Alignment>
58 class RegularTileAccessIterator<
59 Shape_, Element_,
60 layout::TensorOpMultiplicandCongruous<sizeof_bits<Element_>::value,
61 int(128 / sizeof(Element_))>,
62 AdvanceRank, ThreadMap_, Alignment> {
63 public:
64 static_assert(
65 AdvanceRank == 0 || AdvanceRank == 1,
66 "Specialization for pitch-linear iterator may along advance along the "
67 "contiguous(rank=0) or strided(rank=1) dimension.");
68
69 using Shape = Shape_;
70 using Element = Element_;
71 using Layout =
72 layout::TensorOpMultiplicandCongruous<sizeof_bits<Element_>::value,
73 int(128 / sizeof(Element_))>;
74 static int const kAdvanceRank = AdvanceRank;
75 static int const kAlignment = Alignment;
76
77 using Index = typename Layout::Index;
78 using LongIndex = typename Layout::LongIndex;
79
80 using TensorRef = TensorRef<Element, Layout>;
81 using TensorCoord = typename Layout::TensorCoord;
82
83 using ThreadMap = ThreadMap_;
84
86 struct Detail {
89 static int const kAccessSizeInBits = 128;
90
91 static_assert(sizeof_bits<Element_>::value *
92 ThreadMap::kElementsPerAccess ==
93 kAccessSizeInBits,
94 "This iterator requires a policy whose access size is 128bs");
95
97 static int const kPointerCount =
98 (ThreadMap::Iterations::kStrided > 1 ? 2 : 1);
99 };
100
102 using AccessType = Array<Element, Layout::kElementsPerAccess>;
103
104 private:
105 //
106 // Data members
107 //
108
110 Index stride_;
111
113 AccessType *pointer_[Detail::kPointerCount];
114
116 Index byte_offset_;
117
119 int iteration_contiguous_;
120
122 int iteration_strided_;
123
124 public:
126 CUTLASS_HOST_DEVICE
127 RegularTileAccessIterator(TensorRef ref,
128 int thread_id
129 )
130 : stride_(ref.stride(0) / Layout::kElementsPerAccess),
131 byte_offset_(0) {
132 layout::PitchLinearCoord thread_offset_base =
133 ThreadMap::initial_offset(thread_id);
134
135 CUTLASS_PRAGMA_UNROLL
136 for (int i = 0; i < Detail::kPointerCount; ++i) {
137 // This is the offset of a thread within a threadblock tile for a specific
138 // pointer (units of elements)
139 layout::PitchLinearCoord thread_offset_in_threadblock_tile =
140 thread_offset_base +
141 layout::PitchLinearCoord{
142 0, ThreadMap::Detail::WarpThreadArrangement::kStrided * i};
143
144 // initialize pointer
145 pointer_[i] = reinterpret_cast<AccessType *>(
146 ref.data() + ref.offset(thread_offset_in_threadblock_tile));
147 }
148
149 set_iteration_index(0);
150 }
151
153 CUTLASS_HOST_DEVICE
154 void set_iteration_index(int index) {
155 iteration_contiguous_ = index % ThreadMap::Iterations::kContiguous;
156 iteration_strided_ = index / ThreadMap::Iterations::kContiguous;
157 }
158
160 CUTLASS_HOST_DEVICE
161 void add_pointer_offset(LongIndex pointer_offset) {
162 byte_offset_ += pointer_offset * sizeof(Element);
163 }
164
166 CUTLASS_HOST_DEVICE
167 AccessType *get() const {
168 AccessType *access_ptr = pointer_[iteration_strided_ & 1];
169 int stride_idx = (iteration_strided_ & ~1);
170
171 int access_offset = stride_idx * ThreadMap::Delta::kStrided * stride_ +
172 iteration_contiguous_ * ThreadMap::Delta::kContiguous /
173 ThreadMap::kElementsPerAccess;
174
175 char *access_byte_ptr =
176 reinterpret_cast<char *>(access_ptr + access_offset);
177 return reinterpret_cast<AccessType *>(access_byte_ptr + byte_offset_);
178 }
179
181 CUTLASS_HOST_DEVICE
182 RegularTileAccessIterator &operator++() {
183 ++iteration_contiguous_;
184
185 if (iteration_contiguous_ < ThreadMap::Iterations::kContiguous)
186 return *this;
187
188 // Enter here only if (iteration_contiguous_ ==
189 // ThreadMap::Iteration::kContiguous)
190 iteration_contiguous_ = 0;
191 ++iteration_strided_;
192
193 if (iteration_strided_ < ThreadMap::Iterations::kStrided) {
194 return *this;
195 }
196
197 // Enter here only if (iteration_strided_ == ThreadMap::Iteration::kStrided)
198 // which means we enter the next tile.
199 iteration_strided_ = 0;
200
201 return *this;
202 }
203
205 CUTLASS_HOST_DEVICE
206 RegularTileAccessIterator operator++(int) {
207 RegularTileAccessIterator prev(*this);
208 this->operator++();
209
210 return prev;
211 }
212
214 CUTLASS_DEVICE
215 void add_tile_offset(TensorCoord const &coord) {
216 add_pointer_offset(coord.contiguous() * Shape::kContiguous +
217 coord.strided() * Shape::kStrided * stride_ *
218 Layout::kElementsPerAccess);
219 }
220 };
221
223
231 template <typename Shape_, typename Element_, int AdvanceRank,
232 typename ThreadMap_, int Alignment>
233 class RegularTileAccessIterator<
234 Shape_, Element_,
235 layout::ColumnMajorTensorOpMultiplicandCongruous<
236 sizeof_bits<Element_>::value, int(128 / sizeof(Element_))>,
237 AdvanceRank, ThreadMap_, Alignment> {
238 public:
239 static_assert(
240 AdvanceRank == 0 || AdvanceRank == 1,
241 "Specialization for column-major iterator may along advance along the "
242 "columns(rank=0) or rows(rank=1) dimension.");
243
244 using Shape = Shape_;
245 using Element = Element_;
246 using Layout = layout::ColumnMajorTensorOpMultiplicandCongruous<
247 sizeof_bits<Element_>::value, int(128 / sizeof(Element_))>;
248 static int const kAdvanceRank = AdvanceRank;
249 static int const kAlignment = Alignment;
250
251 using Index = typename Layout::Index;
252 using LongIndex = typename Layout::LongIndex;
253
254 using TensorRef = TensorRef<Element, Layout>;
255 using TensorCoord = typename Layout::TensorCoord;
256
257 using ThreadMap = ThreadMap_;
258
260 using UnderlyingIterator = RegularTileAccessIterator<
261 layout::PitchLinearShape<Shape::kRow, Shape::kColumn>, Element,
262 layout::TensorOpMultiplicandCongruous<sizeof_bits<Element_>::value,
263 int(128 / sizeof(Element_))>,
264 (kAdvanceRank == 0 ? 0 : 1), ThreadMap_>;
265
266 using AccessType = typename UnderlyingIterator::AccessType;
267
268 private:
270 UnderlyingIterator iterator_;
271
272 public:
274 CUTLASS_HOST_DEVICE
275 RegularTileAccessIterator(TensorRef ref,
276 int thread_id
277 )
278 : iterator_({ref.data(), ref.stride()}, thread_id) {}
279
281 CUTLASS_HOST_DEVICE
282 void set_iteration_index(int index) { iterator_.set_iteration_index(index); }
283
285 CUTLASS_HOST_DEVICE
286 void add_pointer_offset(LongIndex pointer_offset) {
287 iterator_.add_pointer_offset(pointer_offset);
288 }
289
291 CUTLASS_HOST_DEVICE
292 AccessType *get() const {
293 return reinterpret_cast<AccessType *>(iterator_.get());
294 }
295
297 CUTLASS_DEVICE
298 void add_tile_offset(TensorCoord const &coord) {
299 iterator_.add_tile_offset({coord.row(), coord.column()});
300 }
301
303 CUTLASS_HOST_DEVICE
304 RegularTileAccessIterator &operator++() {
305 ++iterator_;
306 return *this;
307 }
308
310 CUTLASS_HOST_DEVICE
311 RegularTileAccessIterator operator++(int) {
312 RegularTileAccessIterator prev(*this);
313 ++iterator_;
314
315 return prev;
316 }
317 };
318
320
328 template <typename Shape_, typename Element_, int AdvanceRank,
329 typename ThreadMap_, int Alignment>
330 class RegularTileAccessIterator<
331 Shape_, Element_,
332 layout::RowMajorTensorOpMultiplicandCongruous<sizeof_bits<Element_>::value,
333 int(128 / sizeof(Element_))>,
334 AdvanceRank, ThreadMap_, Alignment> {
335 public:
336 static_assert(
337 AdvanceRank == 0 || AdvanceRank == 1,
338 "Specialization for row-major iterator may along advance along the "
339 "columns(rank=0) or rows(rank=1) dimension.");
340
341 using Shape = Shape_;
342 using Element = Element_;
343 using Layout = layout::RowMajorTensorOpMultiplicandCongruous<
344 sizeof_bits<Element_>::value, int(128 / sizeof(Element_))>;
345 static int const kAdvanceRank = AdvanceRank;
346 static int const kAlignment = Alignment;
347
348 using Index = typename Layout::Index;
349 using LongIndex = typename Layout::LongIndex;
350
351 using TensorRef = TensorRef<Element, Layout>;
352 using TensorCoord = typename Layout::TensorCoord;
353
354 using ThreadMap = ThreadMap_;
355
357 using UnderlyingIterator = RegularTileAccessIterator<
358 layout::PitchLinearShape<Shape::kColumn, Shape::kRow>, Element,
359 layout::TensorOpMultiplicandCongruous<sizeof_bits<Element_>::value,
360 int(128 / sizeof(Element_))>,
361 (kAdvanceRank == 0 ? 1 : 0), ThreadMap_>;
362
363 using AccessType = typename UnderlyingIterator::AccessType;
364
365 private:
367 UnderlyingIterator iterator_;
368
369 public:
371 CUTLASS_HOST_DEVICE
372 RegularTileAccessIterator(TensorRef ref,
373 int thread_id
374 )
375 : iterator_({ref.data(), ref.stride()}, thread_id) {}
376
378 CUTLASS_HOST_DEVICE
379 void set_iteration_index(int index) { iterator_.set_iteration_index(index); }
380
382 CUTLASS_HOST_DEVICE
383 void add_pointer_offset(LongIndex pointer_offset) {
384 iterator_.add_pointer_offset(pointer_offset);
385 }
386
388 CUTLASS_HOST_DEVICE
389 AccessType *get() const {
390 return reinterpret_cast<AccessType *>(iterator_.get());
391 }
392
394 CUTLASS_DEVICE
395 void add_tile_offset(TensorCoord const &coord) {
396 iterator_.add_tile_offset({coord.column(), coord.row()});
397 }
398
400 CUTLASS_HOST_DEVICE
401 RegularTileAccessIterator &operator++() {
402 ++iterator_;
403 return *this;
404 }
405
407 CUTLASS_HOST_DEVICE
408 RegularTileAccessIterator operator++(int) {
409 RegularTileAccessIterator prev(*this);
410 ++iterator_;
411
412 return prev;
413 }
414 };
415
417
425 template <typename Shape_, typename Element_, int AdvanceRank,
426 typename ThreadMap_, int Alignment, int Crosswise>
427 class RegularTileAccessIterator<Shape_, Element_,
428 layout::TensorOpMultiplicandCrosswise<
429 sizeof_bits<Element_>::value, Crosswise>,
430 AdvanceRank, ThreadMap_, Alignment> {
431 public:
432 static_assert(
433 AdvanceRank == 0 || AdvanceRank == 1,
434 "Specialization for pitch-linear iterator may along advance along the "
435 "contiguous(rank=0) or strided(rank=1) dimension.");
436
437 using Shape = Shape_;
438 using Element = Element_;
439 using Layout =
440 layout::TensorOpMultiplicandCrosswise<sizeof_bits<Element_>::value,
441 Crosswise>;
442 static int const kAdvanceRank = AdvanceRank;
443 static int const kAlignment = Alignment;
444 static int const kCrosswise = Crosswise;
445
446 using Index = typename Layout::Index;
447 using LongIndex = typename Layout::LongIndex;
448
449 using TensorRef = TensorRef<Element, Layout>;
450 using TensorCoord = typename Layout::TensorCoord;
451
452 using ThreadMap = ThreadMap_;
453
455 struct Detail {
458 static int const kAccessSizeInBits = 128;
459
460 static_assert(sizeof_bits<Element_>::value *
461 ThreadMap::kElementsPerAccess ==
462 kAccessSizeInBits,
463 "This iterator requires a policy whose access size is 128bs");
464
469 static int const kPointerCount =
470 (ThreadMap::Iterations::kStrided > 1 ? 2 : 1);
471 };
472
474 using AccessType = Array<Element, Layout::kElementsPerAccess>;
475
476 private:
477 //
478 // Data members
479 //
480
485 int sections_;
486
488 int sections_per_stage_;
489
491 Index stride_;
492
494 AccessType *pointer_[Detail::kPointerCount];
495
497 Index byte_offset_;
498
500 int iteration_contiguous_;
501
503 int iteration_strided_;
504
505 public:
507 CUTLASS_HOST_DEVICE
508 RegularTileAccessIterator(TensorRef ref,
509 int thread_id
510 )
511 : sections_(ref.stride(0) / kCrosswise),
512 sections_per_stage_(Shape::kContiguous / kCrosswise),
513 // stride_ = kCrosswise x sections_ x kFactor
514 stride_(ref.stride(0) * Layout::kFactor / Layout::kElementsPerAccess),
515 byte_offset_(0) {
516 layout::PitchLinearCoord thread_offset_base =
517 ThreadMap::initial_offset(thread_id);
518
519 CUTLASS_PRAGMA_UNROLL
520 for (int i = 0; i < Detail::kPointerCount; ++i) {
521 // This is the offset of a thread within a threadblock tile for a specific
522 // pointer (units of elements)
523 layout::PitchLinearCoord thread_offset_in_threadblock_tile =
524 thread_offset_base +
525 layout::PitchLinearCoord{
526 0, ThreadMap::Detail::WarpThreadArrangement::kStrided * i};
527 // initialize pointer
528 pointer_[i] = reinterpret_cast<AccessType *>(ref.data()) +
529 ref.offset(thread_offset_in_threadblock_tile) /
530 Layout::kElementsPerAccess;
531 }
532
533 set_iteration_index(0);
534 }
535
537 CUTLASS_HOST_DEVICE
538 void set_iteration_index(int index) {
539 iteration_contiguous_ = index % ThreadMap::Iterations::kContiguous;
540 iteration_strided_ = index / ThreadMap::Iterations::kContiguous;
541 }
542
544 CUTLASS_HOST_DEVICE
545 void add_pointer_offset(LongIndex pointer_offset) {
546 byte_offset_ += pointer_offset * sizeof_bits<Element>::value / 8;
547 }
548
550 CUTLASS_HOST_DEVICE
551 AccessType *get() const {
552 AccessType *access_ptr = pointer_[iteration_strided_ & 1];
553 int stride_idx = (iteration_strided_ & ~1);
554
555 int access_offset =
556 stride_idx * ThreadMap::Delta::kStrided * stride_ / Layout::kFactor +
557 iteration_contiguous_ * ThreadMap::Delta::kContiguous /
558 ThreadMap::kElementsPerAccess;
559 char *access_byte_ptr =
560 reinterpret_cast<char *>(access_ptr + access_offset);
561 return reinterpret_cast<AccessType *>(access_byte_ptr + byte_offset_);
562 }
563
565 CUTLASS_HOST_DEVICE
566 RegularTileAccessIterator &operator++() {
567 ++iteration_contiguous_;
568
569 if (iteration_contiguous_ < ThreadMap::Iterations::kContiguous)
570 return *this;
571
572 // Enter here only if (iteration_contiguous_ ==
573 // ThreadMap::Iteration::kContiguous)
574 iteration_contiguous_ = 0;
575 ++iteration_strided_;
576
577 if (iteration_strided_ < ThreadMap::Iterations::kStrided) {
578 return *this;
579 }
580
581 // Enter here only if (iteration_strided_ == ThreadMap::Iteration::kStrided)
582 // which means we enter the next section.
583 iteration_strided_ = 0;
584
585 return *this;
586 }
587
589 CUTLASS_HOST_DEVICE
590 RegularTileAccessIterator operator++(int) {
591 RegularTileAccessIterator prev(*this);
592 this->operator++();
593
594 return prev;
595 }
596
598 CUTLASS_DEVICE
599 void add_tile_offset(TensorCoord const &coord) {
600 add_pointer_offset(coord.contiguous() * sections_per_stage_ * stride_ *
601 ThreadMap::kElementsPerAccess / sections_ +
602 coord.strided() * Shape::kStrided * stride_ *
603 Layout::kElementsPerAccess);
604 }
605 };
606
608
616 template <typename Shape_, typename Element_, int AdvanceRank,
617 typename ThreadMap_, int Alignment, int Crosswise>
618 class RegularTileAccessIterator<
619 Shape_, Element_,
620 layout::ColumnMajorTensorOpMultiplicandCrosswise<
621 sizeof_bits<Element_>::value, Crosswise>,
622 AdvanceRank, ThreadMap_, Alignment> {
623 public:
624 static_assert(
625 AdvanceRank == 0 || AdvanceRank == 1,
626 "Specialization for column-major iterator may along advance along the "
627 "columns(rank=0) or rows(rank=1) dimension.");
628
629 using Shape = Shape_;
630 using Element = Element_;
631 using Layout = layout::ColumnMajorTensorOpMultiplicandCrosswise<
632 sizeof_bits<Element_>::value, Crosswise>;
633 static int const kAdvanceRank = AdvanceRank;
634 static int const kAlignment = Alignment;
635
636 using Index = typename Layout::Index;
637 using LongIndex = typename Layout::LongIndex;
638
639 using TensorRef = TensorRef<Element, Layout>;
640 using TensorCoord = typename Layout::TensorCoord;
641
642 using ThreadMap = ThreadMap_;
643
645 using UnderlyingIterator = RegularTileAccessIterator<
646 layout::PitchLinearShape<Shape::kRow, Shape::kColumn>, Element,
647 layout::TensorOpMultiplicandCrosswise<sizeof_bits<Element_>::value,
648 Crosswise>,
649 (kAdvanceRank == 0 ? 0 : 1), ThreadMap_>;
650
651 using AccessType = typename UnderlyingIterator::AccessType;
652
653 private:
655 UnderlyingIterator iterator_;
656
657 public:
659 CUTLASS_HOST_DEVICE
660 RegularTileAccessIterator(TensorRef ref,
661 int thread_id
662 )
663 : iterator_({ref.data(), ref.stride()}, thread_id) {}
664
666 CUTLASS_HOST_DEVICE
667 void set_iteration_index(int index) { iterator_.set_iteration_index(index); }
668
670 CUTLASS_HOST_DEVICE
671 void add_pointer_offset(LongIndex pointer_offset) {
672 iterator_.add_pointer_offset(pointer_offset);
673 }
674
676 CUTLASS_HOST_DEVICE
677 AccessType *get() const {
678 return reinterpret_cast<AccessType *>(iterator_.get());
679 }
680
682 CUTLASS_DEVICE
683 void add_tile_offset(TensorCoord const &coord) {
684 iterator_.add_tile_offset({coord.row(), coord.column()});
685 }
686
688 CUTLASS_HOST_DEVICE
689 RegularTileAccessIterator &operator++() {
690 ++iterator_;
691 return *this;
692 }
693
695 CUTLASS_HOST_DEVICE
696 RegularTileAccessIterator operator++(int) {
697 RegularTileAccessIterator prev(*this);
698 ++iterator_;
699
700 return prev;
701 }
702 };
703
705
713 template <typename Shape_, typename Element_, int AdvanceRank,
714 typename ThreadMap_, int Alignment, int Crosswise>
715 class RegularTileAccessIterator<Shape_, Element_,
716 layout::RowMajorTensorOpMultiplicandCrosswise<
717 sizeof_bits<Element_>::value, Crosswise>,
718 AdvanceRank, ThreadMap_, Alignment> {
719 public:
720 static_assert(
721 AdvanceRank == 0 || AdvanceRank == 1,
722 "Specialization for row-major iterator may along advance along the "
723 "columns(rank=0) or rows(rank=1) dimension.");
724
725 using Shape = Shape_;
726 using Element = Element_;
727 using Layout = layout::RowMajorTensorOpMultiplicandCrosswise<
728 sizeof_bits<Element_>::value, Crosswise>;
729 static int const kAdvanceRank = AdvanceRank;
730 static int const kAlignment = Alignment;
731
732 using Index = typename Layout::Index;
733 using LongIndex = typename Layout::LongIndex;
734
735 using TensorRef = TensorRef<Element, Layout>;
736 using TensorCoord = typename Layout::TensorCoord;
737
738 using ThreadMap = ThreadMap_;
739
741 using UnderlyingIterator = RegularTileAccessIterator<
742 layout::PitchLinearShape<Shape::kColumn, Shape::kRow>, Element,
743 layout::TensorOpMultiplicandCrosswise<sizeof_bits<Element_>::value,
744 Crosswise>,
745 (kAdvanceRank == 0 ? 1 : 0), ThreadMap_>;
746
747 using AccessType = typename UnderlyingIterator::AccessType;
748
749 private:
751 UnderlyingIterator iterator_;
752
753 public:
755 CUTLASS_HOST_DEVICE
756 RegularTileAccessIterator(TensorRef ref,
757 int thread_id
758 )
759 : iterator_({ref.data(), ref.stride()}, thread_id) {}
760
762 CUTLASS_HOST_DEVICE
763 void set_iteration_index(int index) { iterator_.set_iteration_index(index); }
764
766 CUTLASS_HOST_DEVICE
767 void add_pointer_offset(LongIndex pointer_offset) {
768 iterator_.add_pointer_offset(pointer_offset);
769 }
770
772 CUTLASS_HOST_DEVICE
773 AccessType *get() const {
774 return reinterpret_cast<AccessType *>(iterator_.get());
775 }
776
778 CUTLASS_DEVICE
779 void add_tile_offset(TensorCoord const &coord) {
780 iterator_.add_tile_offset({coord.column(), coord.row()});
781 }
782
784 CUTLASS_HOST_DEVICE
785 RegularTileAccessIterator &operator++() {
786 ++iterator_;
787 return *this;
788 }
789
791 CUTLASS_HOST_DEVICE
792 RegularTileAccessIterator operator++(int) {
793 RegularTileAccessIterator prev(*this);
794 ++iterator_;
795
796 return prev;
797 }
798 };
799
801 } // namespace threadblock
802 } // namespace transform
803 } // namespace cutlass
804
cutlass::layout::ColumnMajorTensorOpMultiplicandCongruous::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:434
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:683
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:508
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:640
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:275
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:395
cutlass::layout::TensorOpMultiplicandCrosswise::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:640
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:251
cutlass
Definition: aligned_buffer.h:35
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:538
cutlass::layout::PitchLinearCoord
Coordinate in pitch-linear space.
Definition: pitch_linear.h:52
tensor_ref.h
Defines a structure containing strides, bounds, and a pointer to tensor data.
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:566
cutlass::layout::ColumnMajorTensorOpMultiplicandCongruous::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:431
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:725
tensor_op_multiplicand_sm75.h
cutlass::layout::RowMajorTensorOpMultiplicandCongruous::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:536
cutlass::layout::ColumnMajorTensorOpMultiplicandCrosswise
Definition: tensor_op_multiplicand_sm75.h:734
cutlass::TensorRef::data
CUTLASS_HOST_DEVICE Element * data() const
Returns the pointer to referenced data.
Definition: tensor_ref.h:254
cutlass::layout::RowMajorTensorOpMultiplicandCongruous::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:539
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:629
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:372
cutlass::layout::TensorOpMultiplicandCongruous::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:221
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:127
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:341
cutlass::layout::ColumnMajorTensorOpMultiplicandCongruous
Definition: tensor_op_multiplicand_sm75.h:422
cutlass::layout::RowMajorTensorOpMultiplicandCrosswise::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:843
cutlass::layout::ColumnMajorTensorOpMultiplicandCrosswise::Index
int32_t Index
Index type used for coordinates.
Definition: tensor_op_multiplicand_sm75.h:742
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:642
cutlass::layout::RowMajorTensorOpMultiplicandCrosswise
Definition: tensor_op_multiplicand_sm75.h:835
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:401
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:77
cutlass::layout::TensorOpMultiplicandCongruous
Definition: tensor_op_multiplicand_sm75.h:213
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:298
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:590
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:154
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:452
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:545
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:792
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:282
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:354
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:161
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::AccessType
typename UnderlyingIterator::AccessType AccessType
Definition: regular_tile_access_iterator_tensor_op.h:747
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:437
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:446
cutlass::layout::TensorOpMultiplicandCongruous::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:224
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:637
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:408
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:81
cutlass::layout::PitchLinearShape
Template defining a shape used by pitch-linear operators.
Definition: pitch_linear.h:43
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::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:767
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:450
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:352
cutlass::operator++
CUTLASS_HOST_DEVICE half_t & operator++(half_t &lhs)
Definition: half.h:694
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:78
cutlass::layout::RowMajorTensorOpMultiplicandCrosswise::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:846
cutlass::TensorRef::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the layout object&#39;s stride vector.
Definition: tensor_ref.h:277
matrix_shape.h
Defines a Shape template for matrix tiles.
cutlass::layout::ColumnMajorTensorOpMultiplicandCrosswise::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:745
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:348
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:255
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:206
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:689
cutlass::TensorRef< Element, Layout >
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:379
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:636
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Element
Element_ Element
Definition: regular_tile_access_iterator_tensor_op.h:245
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:182
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Element
Element_ Element
Definition: regular_tile_access_iterator_tensor_op.h:70
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::Index
typename Layout::Index Index
Definition: regular_tile_access_iterator_tensor_op.h:732
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:349
cutlass::TensorRef::offset
CUTLASS_HOST_DEVICE LongIndex offset(TensorCoord const &coord) const
Computes the offset of an index from the origin of the tensor.
Definition: tensor_ref.h:301
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::AccessType
Array< Element, Layout::kElementsPerAccess > AccessType
Element type per access.
Definition: regular_tile_access_iterator_tensor_op.h:102
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:763
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:738
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:599
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:671
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:304
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:244
cutlass::layout::TensorOpMultiplicandCrosswise::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: tensor_op_multiplicand_sm75.h:643
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::AccessType
Array< Element, Layout::kElementsPerAccess > AccessType
Element type per access.
Definition: regular_tile_access_iterator_tensor_op.h:474
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:696
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:383
regular_tile_access_iterator.h
Templates implementing the address computation of storing of tiles from pitch-linear rank=2 tensors...
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:660
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::AccessType
typename UnderlyingIterator::AccessType AccessType
Definition: regular_tile_access_iterator_tensor_op.h:651
matrix_coord.h
Defines a canonical coordinate for rank=2 matrices offering named indices.
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator operator++(int)
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:311
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex pointer_offset)
Adds a pointer offset in units of Element.
Definition: regular_tile_access_iterator_tensor_op.h:286
cutlass::transform::threadblock::RegularTileAccessIterator
Definition: regular_tile_access_iterator.h:48
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::AccessType
typename UnderlyingIterator::AccessType AccessType
Definition: regular_tile_access_iterator_tensor_op.h:363
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:257
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::AccessType
typename UnderlyingIterator::AccessType AccessType
Definition: regular_tile_access_iterator_tensor_op.h:266
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Shape
Shape_ Shape
Definition: regular_tile_access_iterator_tensor_op.h:69
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::set_iteration_index
CUTLASS_HOST_DEVICE void set_iteration_index(int index)
Overrides the internal iteration index.
Definition: regular_tile_access_iterator_tensor_op.h:667
pitch_linear.h
Defines layout functions used by TensorRef and derived classes for pitch-linear memory.
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::operator++
CUTLASS_HOST_DEVICE RegularTileAccessIterator & operator++()
Advances to the next tile in memory.
Definition: regular_tile_access_iterator_tensor_op.h:785
cutlass::layout::TensorOpMultiplicandCrosswise
Definition: tensor_op_multiplicand_sm75.h:632
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::TensorCoord
typename Layout::TensorCoord TensorCoord
Definition: regular_tile_access_iterator_tensor_op.h:736
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:779
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:252
cutlass.h
Basic include for CUTLASS.
cutlass::MatrixCoord
Definition: matrix_coord.h:39
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:733
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::LongIndex
typename Layout::LongIndex LongIndex
Definition: regular_tile_access_iterator_tensor_op.h:447
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Crosswise >, AdvanceRank, ThreadMap_, Alignment >::RegularTileAccessIterator
CUTLASS_HOST_DEVICE RegularTileAccessIterator(TensorRef ref, int thread_id)
Construct a TileIterator with zero threadblock offset.
Definition: regular_tile_access_iterator_tensor_op.h:756
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::ThreadMap
ThreadMap_ ThreadMap
Definition: regular_tile_access_iterator_tensor_op.h:83
cutlass::transform::threadblock::RegularTileAccessIterator< Shape_, Element_, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset
CUTLASS_DEVICE void add_tile_offset(TensorCoord const &coord)
Adds a tile offset.
Definition: regular_tile_access_iterator_tensor_op.h:215
cutlass::layout::RowMajorTensorOpMultiplicandCongruous
Definition: tensor_op_multiplicand_sm75.h:527