Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

tile_iterator.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 "cutlass/coord.h"
32 #include "cutlass/tensor_ref.h"
33 #include "cutlass/fragment.h"
34 #include "cutlass/load_store.h"
35 #include "cutlass/predicate_vector.h"
36 #include "cutlass/vector.h"
37 #include <cstdio>
38
39 namespace cutlass {
40
42
61
64 struct IteratorAdvance {
65 enum Kind { kD, kH, kW };
66 };
68
73 template <typename Tile_,
74 typename Delta_,
75 typename Iterations_,
76 typename ThreadOffset_,
77 int AccessSize>
78 struct TileTraits {
80 typedef Tile_ Tile;
81
83 typedef Delta_ Delta;
84
86 typedef Iterations_ Iterations;
87
89 //
90 // ThreadOffset should be a functor defined like:
91 //
92 // struct ThreadOffsetExample {
93 // CUTLASS_DEVICE
94 // Coord<4> operator()() const {
95 // return make_Coord(0, threadIdx.y, threadIdx.x, 0);
96 // }
97 // };
98 //
99 typedef ThreadOffset_ ThreadOffset;
100
102 typedef Shape<0, 0, 0, 0> ImmediateOffsetStrides;
103
105 static int const kAccessSize = AccessSize;
106 };
107
109
111 template <typename Delta_>
112 struct RegularTilePredicateFunctor {
113 typedef Delta_ Delta;
114
116 Coord<3> bounds;
117
119 CUTLASS_HOST_DEVICE
120 RegularTilePredicateFunctor(Coord<3> _bounds) : bounds(_bounds) {}
121
123 CUTLASS_HOST_DEVICE
124 bool operator()(Coord<3> iteration, Coord<3> offset) const {
125 return (iteration[0] * Delta::kD + offset[0] < bounds[0]) &&
126 (iteration[1] * Delta::kH + offset[1] < bounds[1]) &&
127 (iteration[2] * Delta::kW + offset[2] < bounds[2]);
128 }
129 };
130
132
133 template <typename T>
134 struct DumpType {};
136 template <typename Traits_,
137 typename Scalar_,
138 IteratorAdvance::Kind Advance_ = IteratorAdvance::kH,
139 MemorySpace::Kind MemorySpace = MemorySpace::kGeneric,
140 typename Index_ = int,
141 typename FragmentElement_ = Scalar_,
142 FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar,
143 typename Skew_ = Shape<0, 0, 0, 0> >
144 struct TileIteratorBase {
146 typedef Traits_ Traits;
147
149 typedef Scalar_ Scalar;
150
152 typedef FragmentElement_ FragmentElement;
153
155 static IteratorAdvance::Kind const kAdvance = Advance_;
156
158 static FragmentElementType::Kind const kFragmentElementType = FragmentElementType_;
159
161 static MemorySpace::Kind const kMemorySpace = MemorySpace;
162
164 typedef Index_ Index;
165
167 typedef long long LongIndex;
168
170 typedef Skew_ Skew;
171
173 typedef typename Traits::Tile Tile;
174
176 typedef typename Traits::Delta Delta;
177
179 typedef typename Traits::ImmediateOffsetStrides ImmediateOffsetStrides;
180
182 typedef typename Traits::Iterations Iterations;
183
185 typedef typename Traits::ThreadOffset ThreadOffset;
186
188 static int const kAccessSize = Traits::kAccessSize;
189
191 typedef typename Vectorize<FragmentElement, kAccessSize>::Type AccessType;
192
194 static int const kFragmentSize =
195 (kFragmentElementType == FragmentElementType::kWmmaMatrix ? 16 : sizeof(AccessType));
197 typedef Fragment<Scalar, ShapeCount<Tile>::kCount, kFragmentSize> Storage;
199 typedef Fragment<FragmentElement, ShapeCount<Iterations>::kCount * kAccessSize> Fragment;
200
202 typedef FragmentIterator<Fragment, Iterations, AccessType> FragmentIterator;
204 typedef FragmentConstIterator<Fragment, Iterations, AccessType> FragmentConstIterator;
206 typedef typename FragmentIterator::FragmentShape FragmentShape;
207
209 typedef PredicateVector<ShapeCount<Iterations>::kCount> PredicateVector;
210
211 //
212 // Params struct
213 //
214
216 struct Params {
217
218 //
219 // Dat members
220 //
221
222 Index stride_d;
223 Index stride_h;
224 Index stride_w;
225
226 Index inc_d;
227 Index inc_h;
228 Index inc_w;
229
230 Index inc_advance;
231
232 //
233 // Methods
234 //
235
237 CUTLASS_HOST_DEVICE
238 Params() : stride_d(0), stride_h(0), stride_w(0), inc_d(0), inc_h(0), inc_w(0) {}
239
241 CUTLASS_HOST_DEVICE
242 Params(Index _stride_d,
243 Index _stride_h,
244 Index _stride_w,
245 Index _inc_d,
246 Index _inc_h,
247 Index _inc_w,
248 Index _inc_advance)
249 : stride_d(_stride_d),
250 stride_h(_stride_h),
251 stride_w(_stride_w),
252 inc_d(_inc_d),
253 inc_h(_inc_h),
254 inc_w(_inc_w),
255 inc_advance(_inc_advance) {}
256
258 CUTLASS_HOST_DEVICE
259 Params(Coord<4> const &stride) {
260 initialize(stride);
261 }
262
264 CUTLASS_HOST_DEVICE
265 int initialize(Index _stride_d,
266 Index _stride_h,
267 Index _stride_w,
268 Index _inc_d,
269 Index _inc_h,
270 Index _inc_w,
271 Index _inc_advance) {
272 stride_d = _stride_d;
273 stride_h = _stride_h;
274 stride_w = _stride_w;
275
276 inc_d = _inc_d;
277 inc_h = _inc_h;
278 inc_w = _inc_w;
279 inc_advance = _inc_advance;
280
281 return 0;
282 }
283
285 CUTLASS_HOST_DEVICE
286 int initialize(Coord<4> const &stride) {
287 return initialize(stride[0], stride[1], stride[2]);
288 }
289
291 CUTLASS_HOST_DEVICE
292 int initialize(Index _stride_d, Index _stride_h, Index _stride_w) {
293 stride_d = _stride_d;
294 stride_h = _stride_h;
295 stride_w = _stride_w;
296
297 inc_w = stride_w * Delta::kW;
298 inc_h = stride_h * Delta::kH - stride_w * Delta::kW * (Iterations::kW - 1);
299 inc_d = stride_h * Delta::kD - stride_h * Delta::kH * (Iterations::kH - 1) -
300 stride_w * Delta::kW * (Iterations::kW - 1);
301
302 inc_advance = 0;
303
304 if (kAdvance == IteratorAdvance::kH) {
305 // Advance in the H dimension.
306 inc_advance = Tile::kH * stride_h;
307 } else if (kAdvance == IteratorAdvance::kW) {
308 // Advance in the W dimension.
309 inc_advance = Tile::kW * stride_w;
310
311 } else {
312 // Advance in the D dimension.
313 inc_advance = Tile::kD * stride_d;
314 }
315
316 inc_advance -= stride_h * Delta::kD * (Iterations::kD - 1) +
317 stride_h * Delta::kH * (Iterations::kH - 1) +
318 stride_w * Delta::kW * (Iterations::kW - 1);
319
320 return 0;
321 }
322
324 CUTLASS_HOST_DEVICE int initialize() {
325 stride_d = 0;
326 stride_h = 0;
327 stride_w = 1;
328
329 inc_advance = 0;
330 inc_d = inc_h = inc_w = 0;
331
332 return 0;
333 }
334 };
335
337 CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const { return true; }
338
339 //
340 // Static function members
341 //
342
344 template <typename PredicateIterator, typename PredicateFunctor>
345 CUTLASS_HOST_DEVICE static void initialize_predicates(PredicateIterator predicate_it,
346 PredicateFunctor const &predicate_func,
347 Coord<3> const &offset) {
348 CUTLASS_PRAGMA_UNROLL
349 for (int d = 0; d < Iterations::kD; ++d) {
350 CUTLASS_PRAGMA_UNROLL
351 for (int h = 0; h < Iterations::kH; ++h) {
352 CUTLASS_PRAGMA_UNROLL
353 for (int w = 0; w < Iterations::kW; ++w) {
354 bool enable = predicate_func(make_Coord(d, h, w), offset);
355 predicate_it.set(enable);
356 ++predicate_it;
357 }
358 }
359 }
360 }
361 };
362
364
388
394 template <typename Traits_,
395 typename Scalar_,
396 IteratorAdvance::Kind Advance_ = IteratorAdvance::kH,
397 MemorySpace::Kind MemorySpace = MemorySpace::kGeneric,
398 typename Index_ = int,
399 typename FragmentElement_ = Scalar_,
400 FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar,
401 typename Skew_ = Shape<0, 0, 0, 0> >
402 struct TileLoadIterator : public TileIteratorBase<Traits_,
403 Scalar_,
404 Advance_,
405 MemorySpace,
406 Index_,
407 FragmentElement_,
408 FragmentElementType_,
409 Skew_> {
411 typedef TileIteratorBase<Traits_,
412 Scalar_,
413 Advance_,
414 MemorySpace,
415 Index_,
416 FragmentElement_,
417 FragmentElementType_,
418 Skew_>
419 Base;
420
422 typedef typename Base::Traits Traits;
423
425 typedef typename Base::Scalar Scalar;
426
428 typedef FragmentElement_ FragmentElement;
429
431 static IteratorAdvance::Kind const kAdvance = Base::kAdvance;
432
434 static FragmentElementType::Kind const kFragmentElementType = FragmentElementType_;
435
437 static MemorySpace::Kind const kMemorySpace = Base::kMemorySpace;
438
440 typedef typename Base::Index Index;
441
443 typedef typename Base::LongIndex LongIndex;
444
446 typedef typename Base::Skew Skew;
447
449 typedef typename Base::Tile Tile;
450
452 typedef typename Base::Delta Delta;
453
455 typedef typename Base::Iterations Iterations;
456
458 typedef typename Base::ThreadOffset ThreadOffset;
459
461 typedef typename Base::FragmentShape FragmentShape;
462
464 typedef typename Base::AccessType AccessType;
465
467 static int const kAccessSize = Base::kAccessSize;
468
470 typedef typename Base::Fragment Fragment;
471
473 typedef typename Base::FragmentIterator FragmentIterator;
474
476 typedef typename Base::FragmentConstIterator FragmentConstIterator;
477
479 typedef typename Base::PredicateVector PredicateVector;
480
482 typedef typename Base::Storage SharedStorage;
483
485 typedef typename Base::Params BaseParams;
486
488 enum { kRequiresLoadFence = Tile::kD == 1 };
489
491 typedef Scalar const *Pointer;
492
494 typedef TensorRef<Scalar const, 4> TensorRef;
495
497 struct Params : public BaseParams {
499 Scalar const *pointer;
500
501 //
502 // Methods
503 //
504
506 CUTLASS_HOST_DEVICE
507 Params() : pointer(0){ Base::Params::initialize(); }
508
510 CUTLASS_HOST_DEVICE
511 Params(Scalar const *ptr) : pointer(ptr) { Base::Params::initialize(); }
512
514 CUTLASS_HOST_DEVICE
515 Params(TensorRef const &ref): pointer(ref.data()) {
516 Base::Params::initialize(ref.stride());
517 }
518
520 CUTLASS_HOST_DEVICE
521 Params(Scalar const *ptr,
522 Index _stride_d,
523 Index _stride_h,
524 Index _stride_w,
525 Index _inc_d,
526 Index _inc_h,
527 Index _inc_w,
528 Index _inc_advance)
529 : pointer(ptr) {
530 Base::Params::initialize(
531 _stride_d, _stride_h, _stride_w, _inc_d, _inc_h, _inc_w, _inc_advance);
532 }
533
535 CUTLASS_HOST_DEVICE
536 Params(Scalar const *ptr, Index stride_d, Index stride_h, Index stride_w)
537 : pointer(ptr) {
538 Base::Params::initialize(stride_d, stride_h, stride_w);
539 }
540
542 CUTLASS_HOST_DEVICE
543 int initialize(TensorRef const &ref) {
544 pointer = ref.data();
545 return Base::Params::initialize(ref.stride());
546 }
547
549 CUTLASS_HOST_DEVICE
550 int initialize(SharedStorage const &storage) {
551 pointer = &storage[0];
552 Base::Params::initialize();
553 return 0;
554 }
555
557 CUTLASS_HOST_DEVICE
558 int initialize(Scalar const *ptr) {
559 pointer = ptr;
560 Base::Params::initialize();
561 return 0;
562 }
563
565 CUTLASS_HOST_DEVICE
566 int initialize(Scalar const *ptr, Index stride_d, Index stride_h, Index stride_w) {
567 Base::Params::initialize(stride_d, stride_h, stride_w);
568 pointer = ptr;
569 return 0;
570 }
571
573 CUTLASS_HOST_DEVICE
574 int initialize(Scalar const *ptr,
575 Index _stride_d,
576 Index _stride_h,
577 Index _stride_w,
578 Index _inc_d,
579 Index _inc_h,
580 Index _inc_w,
581 Index _inc_advance) {
582 pointer = ptr;
583 Base::Params::initialize(
584 _stride_d, _stride_h, _stride_w, _inc_d, _inc_h, _inc_w, _inc_advance);
585 return 0;
586 }
587
588 // Initializes params to default values
589 CUTLASS_HOST_DEVICE
590 int initialize() { return Base::Params::initialize(); }
591 };
592
593 //
594 // Data members
595 //
596
598 Params params;
599
601 Coord<4> thread_offset;
602
604 int stage;
605
606 //
607 // Predicate initialization
608 //
609
611 template <
613 typename PredicateIterator>
614 CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it,
615 Coord<3> const &bounds,
616 Coord<3> const &block_offset = make_Coord(0,
617 0,
618 0)) {
619 Base::initialize_predicates(
620 predicate_it,
621 RegularTilePredicateFunctor<typename Traits::Delta>(bounds),
622 block_offset + make_Coord(thread_offset[0], thread_offset[1], thread_offset[2]));
623 }
624
626 template <
628 typename PredicateIterator,
630 typename PredicateFunctor>
631 CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it,
632 PredicateFunctor const &functor,
633 Coord<3> const &block_offset) {
634 Base::initialize_predicates(
635 predicate_it,
636 functor,
637 block_offset + make_Coord(thread_offset[0], thread_offset[1], thread_offset[2]));
638 }
639
640 //
641 // Methods
642 //
643
645 CUTLASS_HOST_DEVICE
646 TileLoadIterator() {}
647
649 CUTLASS_HOST_DEVICE
650 TileLoadIterator(Params const &_params,
651 Coord<3> const &block_offset = make_Coord(0, 0, 0),
652 ThreadOffset thread_offset_func = ThreadOffset())
653 : params(_params), stage(0) {
654 thread_offset = thread_offset_func();
655
656 Index pointer_offset = Index((block_offset[0] + thread_offset[0]) * params.stride_d) +
657 Index((block_offset[1] + thread_offset[1]) * params.stride_h) +
658 Index((block_offset[2] + thread_offset[2]) * params.stride_w);
659
660 params.pointer += pointer_offset;
661 }
662
664 CUTLASS_HOST_DEVICE
665 TileLoadIterator(Params const &,
666 Scalar const *ptr,
667 Coord<3> const &block_offset = make_Coord(0, 0, 0),
668 ThreadOffset thread_offset_func = ThreadOffset())
669 : stage(0) {
670 params.pointer = ptr + thread_offset_func()[2];
671
672 params.stride_d = 0;
673 params.stride_h = 0;
674 params.stride_w = 1;
675
676 params.inc_d = params.inc_h = params.inc_w = params.inc_advance = 0;
677 }
678
680 CUTLASS_HOST_DEVICE void inc_d() { params.pointer += params.inc_d; }
681
683 CUTLASS_HOST_DEVICE void inc_h() { params.pointer += params.inc_h; }
684
686 CUTLASS_HOST_DEVICE void inc_w() { params.pointer += params.inc_w; }
687
689 CUTLASS_HOST_DEVICE void inc_advance() { params.pointer += params.inc_advance; }
690
692 CUTLASS_HOST_DEVICE void load_element(AccessType &value, int d, int h, int w, int c) const {
693 int const offset =
694 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(d, h, w, c);
695 Load<Scalar,
696 kAccessSize,
697 kMemorySpace,
698 kFragmentElementType,
699 FragmentElement,
700 Tile::kW,
701 sizeof(FragmentElement) * kAccessSize>::load(value, params.pointer, offset);
702 }
703
705 CUTLASS_HOST_DEVICE void inc_stage() {
706 if (Tile::kD > 1) {
707 int const kStageSize = Tile::kH * Tile::kW * Tile::kC;
708 if (stage == Tile::kD - 1) {
709 params.pointer -= (Tile::kD - 1) * kStageSize;
710 stage = 0;
711 } else {
712 params.pointer += kStageSize;
713 stage = stage + 1;
714 }
715 }
716 }
717
719 CUTLASS_HOST_DEVICE TileLoadIterator & operator+=(Coord<3> const &offset) {
720 long long _offset = offset.template dot<long long>(
721 make_Coord(params.stride_d, params.stride_h, params.stride_w)
722 );
723
724 params.pointer += _offset;
725 return *this;
726 }
727
729 CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset) { params.pointer += offset; }
730
731 CUTLASS_HOST_DEVICE Index stride_advance(void) {
732 Index stride = params.stride_h;
733 if (kAdvance == IteratorAdvance::kW) {
734 stride = params.stride_w;
735 }
736 return stride;
737 }
738
740 template <typename Fragment, typename PredicateIterator>
741 CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment, PredicateIterator pred_it) {
742 FragmentIterator frag_iterator(fragment);
743 for (int d = 0; d < Iterations::kD; ++d) {
744 for (int h = 0; h < Iterations::kH; ++h) {
745 for (int w = 0; w < Iterations::kW; ++w, ++pred_it) {
746 for (int c = 0; c < Iterations::kC; ++c) {
747 if (*pred_it) {
748 load_element(
749 reinterpret_cast<AccessType &>(frag_iterator.at(d, h, w, c)), d, h, w, c);
750 }
751 }
752 if (w < Iterations::kW - 1) {
753 inc_w();
754 }
755 }
756 if (h < Iterations::kH - 1) {
757 inc_h();
758 }
759 }
760 if (d < Iterations::kD - 1) {
761 inc_d();
762 }
763 }
764 inc_advance();
765 }
766
768 template <typename Fragment>
769 CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment) {
770 typename PredicateVector::TrivialIterator pred_it;
771 load_post_increment(fragment, pred_it);
772 }
773
775 template <typename Fragment, typename PredicateIterator>
776 CUTLASS_HOST_DEVICE void load(Fragment &fragment, PredicateIterator pred_it) const {
777 TileLoadIterator _load_it(*this);
778 _load_it.load_post_increment(fragment, pred_it);
779 }
780
782 template <typename Fragment>
783 CUTLASS_HOST_DEVICE void load(Fragment &fragment) const {
784 typename PredicateVector::TrivialIterator pred_it;
785 load(fragment, pred_it);
786 }
787
789 template <typename Fragment>
790 CUTLASS_HOST_DEVICE void load(Fragment &fragment, int d) {
791 FragmentIterator frag_iterator(fragment);
792 for (int h = 0; h < Iterations::kH; ++h) {
793 for (int w = 0; w < Iterations::kW; ++w) {
794 for (int c = 0; c < Iterations::kC; ++c) {
795 load_element(reinterpret_cast<AccessType &>(frag_iterator.at(0, h, w, c)), d, h, w, c);
796 }
797 }
798 }
799 }
800 };
801
803
827
833 template <typename Traits_,
834 typename Scalar_,
835 IteratorAdvance::Kind Advance_ = IteratorAdvance::kH,
836 MemorySpace::Kind MemorySpace = MemorySpace::kGeneric,
837 typename Index_ = int,
838 typename FragmentElement_ = Scalar_,
839 FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar,
840 typename Skew_ = Shape<0, 0, 0, 0> >
841 struct TileStoreIterator : public TileIteratorBase<Traits_,
842 Scalar_,
843 Advance_,
844 MemorySpace,
845 Index_,
846 FragmentElement_,
847 FragmentElementType_,
848 Skew_> {
850 typedef TileIteratorBase<Traits_,
851 Scalar_,
852 Advance_,
853 MemorySpace,
854 Index_,
855 FragmentElement_,
856 FragmentElementType_,
857 Skew_>
858 Base;
859
861 typedef typename Base::Traits Traits;
862
864 typedef typename Base::Scalar Scalar;
865
867 typedef typename Base::FragmentElement FragmentElement;
868
870 static IteratorAdvance::Kind const kAdvance = Base::kAdvance;
871
873 static FragmentElementType::Kind const kFragmentElementType = Base::kFragmentElementType;
874
876 static MemorySpace::Kind const kMemorySpace = Base::kMemorySpace;
877
879 static int const kAccessSize = Base::kAccessSize;
880
882 typedef typename Base::Index Index;
883
885 typedef typename Base::LongIndex LongIndex;
886
888 typedef typename Base::Skew Skew;
889
891 typedef typename Base::Tile Tile;
892
894 typedef typename Base::Delta Delta;
895
897 typedef typename Base::Iterations Iterations;
898
900 typedef typename Base::ThreadOffset ThreadOffset;
901
903 typedef typename Base::FragmentShape FragmentShape;
904
906 typedef typename Base::AccessType AccessType;
907
909 typedef typename Base::Fragment Fragment;
910
912 typedef typename Base::FragmentIterator FragmentIterator;
913
915 typedef typename Base::FragmentConstIterator FragmentConstIterator;
916
918 typedef typename Base::PredicateVector PredicateVector;
919
921 typedef typename Base::Storage SharedStorage;
922
924 typedef typename Base::Params BaseParams;
925
927 typedef Scalar *Pointer;
928
930 typedef TensorRef<Scalar, 4> TensorRef;
931
933 struct Params : public BaseParams {
935 Scalar *pointer;
936
937 //
938 // Methods
939 //
940
941 // Default constructor
942 CUTLASS_HOST_DEVICE
943 Params() : pointer(0) {}
944
945 // Default constructor
946 CUTLASS_HOST_DEVICE
947 Params(Scalar *ptr) : pointer(ptr) { Base::Params::initialize(); }
948
950 CUTLASS_HOST_DEVICE
951 Params(TensorRef const &ref): pointer(ref.data()) {
952 Base::Params::initialize(ref.stride());
953 }
954
955 // Default constructor
956 CUTLASS_HOST_DEVICE
957 Params(Scalar *ptr, long long stride_d, Index stride_h, Index stride_w) {
958 initialize(ptr, stride_d, stride_h, stride_w);
959 }
960
961 // Default constructor
962 CUTLASS_HOST_DEVICE
963 Params(Scalar *ptr,
964 Index _stride_d,
965 Index _stride_h,
966 Index _stride_w,
967 Index _inc_d,
968 Index _inc_h,
969 Index _inc_w,
970 Index _inc_advance) {
971 initialize(ptr, _stride_d, _stride_h, _stride_w, _inc_d, _inc_h, _inc_w, _inc_advance);
972 }
973
975 CUTLASS_HOST_DEVICE
976 int initialize(SharedStorage &storage) {
977 pointer = &storage[0];
978 return Base::Params::initialize();
979 }
980
982 CUTLASS_HOST_DEVICE
983 int initialize(Scalar *ptr) {
984 pointer = ptr;
985 return Base::Params::initialize();
986 }
987
989 CUTLASS_HOST_DEVICE
990 int initialize(Scalar *ptr, Index stride_d, Index stride_h, Index stride_w) {
991 Base::Params::initialize(stride_d, stride_h, stride_w);
992 pointer = ptr;
993 return 0;
994 }
995
997 CUTLASS_HOST_DEVICE
998 int initialize(Scalar *ptr,
999 Index _stride_d,
1000 Index _stride_h,
1001 Index _stride_w,
1002 Index _inc_d,
1003 Index _inc_h,
1004 Index _inc_w,
1005 Index _inc_advance) {
1006 pointer = ptr;
1007 Base::Params::initialize(
1008 _stride_d, _stride_h, _stride_w, _inc_d, _inc_h, _inc_w, _inc_advance);
1009 return 0;
1010 }
1011
1013 CUTLASS_HOST_DEVICE
1014 int initialize() { return Base::Params::initialize(); }
1015 };
1016
1017 //
1018 // Data members
1019 //
1020
1022 Params params;
1023
1025 Coord<4> thread_offset;
1026
1028 int stage;
1029
1030 //
1031 // Predicate initialization
1032 //
1033
1035 template <
1037 typename PredicateIterator>
1038 CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it,
1039 Coord<3> const &bounds,
1040 Coord<3> const &block_offset = make_Coord(0,
1041 0,
1042 0)) {
1043 Base::initialize_predicates(
1044 predicate_it,
1045 RegularTilePredicateFunctor<typename Traits::Delta>(bounds),
1046 block_offset + make_Coord(thread_offset[0], thread_offset[1], thread_offset[2]));
1047 }
1048
1050 template <
1052 typename PredicateIterator,
1054 typename PredicateFunctor>
1055 CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it,
1056 PredicateFunctor const &functor,
1057 Coord<3> const &block_offset) {
1058 Base::initialize_predicates(
1059 predicate_it,
1060 functor,
1061 block_offset + make_Coord(thread_offset[0], thread_offset[1], thread_offset[2]));
1062 }
1063
1064 //
1065 // Methods
1066 //
1067
1069 CUTLASS_HOST_DEVICE
1070 TileStoreIterator() {}
1071
1073 CUTLASS_HOST_DEVICE
1074 TileStoreIterator(Params const &_params,
1075 Coord<3> const &block_offset = make_Coord(0, 0, 0),
1076 ThreadOffset thread_offset_func = ThreadOffset())
1077 : params(_params), stage(0) {
1078 thread_offset = thread_offset_func();
1079
1080 params.pointer += (block_offset[0] + thread_offset[0]) * params.stride_d +
1081 (block_offset[1] + thread_offset[1]) * params.stride_h +
1082 (block_offset[2] + thread_offset[2]) * params.stride_w;
1083 }
1084
1086 CUTLASS_HOST_DEVICE
1087 TileStoreIterator(Params const &, Scalar *ptr, ThreadOffset thread_offset_func = ThreadOffset())
1088 : stage(0) {
1089 params.pointer = ptr + thread_offset_func()[2];
1090 params.stride_d = 0;
1091 params.stride_h = 0;
1092 params.stride_w = 1;
1093
1094 params.inc_d = params.inc_h = params.inc_w = params.inc_advance = 0;
1095 }
1096
1098 CUTLASS_HOST_DEVICE void inc_d() { params.pointer += params.inc_d; }
1099
1101 CUTLASS_HOST_DEVICE void inc_h() { params.pointer += params.inc_h; }
1102
1104 CUTLASS_HOST_DEVICE void inc_w() { params.pointer += params.inc_w; }
1105
1107 CUTLASS_HOST_DEVICE void inc_advance() {}
1108
1110 CUTLASS_HOST_DEVICE void inc_stage() {
1111 if (Tile::kD > 1) {
1112 int const kStageSize = Tile::kH * Tile::kW * Tile::kC;
1113 if (stage == Tile::kD - 1) {
1114 params.pointer -= (Tile::kD - 1) * kStageSize;
1115 stage = 0;
1116 } else {
1117 params.pointer += kStageSize;
1118 stage = stage + 1;
1119 }
1120 }
1121 }
1122
1124 CUTLASS_HOST_DEVICE TileStoreIterator & operator+=(Coord<3> const &offset) {
1125 params.pointer += offset.template dot<long long>(
1126 make_Coord(params.stride_d, params.stride_h, params.stride_w)
1127 );
1128 return *this;
1129 }
1130
1132 CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset) { params.pointer += offset; }
1133
1135 CUTLASS_HOST_DEVICE void store_element(AccessType const &value, int d, int h, int w, int c) {
1136 int const offset =
1137 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(d, h, w, c);
1138 Store<Scalar,
1139 kAccessSize,
1140 kMemorySpace,
1141 kFragmentElementType,
1142 FragmentElement,
1143 Tile::kW,
1144 sizeof(FragmentElement) * kAccessSize>::store(value, params.pointer, offset);
1145 }
1146
1148 template <typename Fragment, typename PredicateIterator>
1149 CUTLASS_HOST_DEVICE void store_post_increment(Fragment const &fragment, PredicateIterator pred_it) {
1150 FragmentConstIterator frag_iterator(fragment);
1151
1152 for (int d = 0; d < Iterations::kD; ++d) {
1153 for (int h = 0; h < Iterations::kH; ++h) {
1154 for (int w = 0; w < Iterations::kW; ++w, ++pred_it) {
1155 for (int c = 0; c < Iterations::kC; ++c) {
1156 if (*pred_it) {
1157 store_element(
1158 reinterpret_cast<AccessType const &>(frag_iterator.at(d, h, w, c)), d, h, w, c);
1159 }
1160 }
1161 if (w < Iterations::kW - 1) {
1162 inc_w();
1163 }
1164 }
1165 if (h < Iterations::kH - 1) {
1166 inc_h();
1167 }
1168 }
1169 if (d < Iterations::kD - 1) {
1170 inc_d();
1171 }
1172 }
1173 inc_advance();
1174 }
1175
1177 template <typename Fragment>
1178 CUTLASS_HOST_DEVICE void store_post_increment(Fragment const &fragment) {
1179 typename PredicateVector::TrivialIterator pred_it;
1180 store_post_increment(fragment, pred_it);
1181 }
1182
1184 template <typename Fragment, typename PredicateIterator>
1185 CUTLASS_HOST_DEVICE void store(Fragment const &fragment, PredicateIterator pred_it) const {
1186 TileStoreIterator _store_it(*this);
1187 _store_it.store_post_increment(fragment, pred_it);
1188 }
1189
1191 template <typename Fragment>
1192 CUTLASS_HOST_DEVICE void store(Fragment const &fragment) const {
1193 typename PredicateVector::TrivialIterator pred_it;
1194 store(fragment, pred_it);
1195 }
1196
1198 CUTLASS_HOST_DEVICE void load_element(AccessType &value, int d, int h, int w, int c) const {
1199 int const offset =
1200 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(d, h, w, c);
1201
1202 Load<Scalar,
1203 kAccessSize,
1204 kMemorySpace,
1205 kFragmentElementType,
1206 FragmentElement,
1207 Tile::kW,
1208 sizeof(FragmentElement) * kAccessSize>::load(value, params.pointer, offset);
1209 }
1210
1212 template <typename Fragment, typename PredicateIterator>
1213 CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment, PredicateIterator pred_it) {
1214 FragmentIterator frag_iterator(fragment);
1215
1216 for (int d = 0; d < Iterations::kD; ++d) {
1217 for (int h = 0; h < Iterations::kH; ++h) {
1218 for (int w = 0; w < Iterations::kW; ++w, ++pred_it) {
1219 for (int c = 0; c < Iterations::kC; ++c) {
1220 if (*pred_it) {
1221 load_element(
1222 reinterpret_cast<AccessType &>(frag_iterator.at(d, h, w, c)), d, h, w, c);
1223 }
1224 }
1225 if (w < Iterations::kW - 1) {
1226 inc_w();
1227 }
1228 }
1229 if (h < Iterations::kH - 1) {
1230 inc_h();
1231 }
1232 }
1233 if (d < Iterations::kD - 1) {
1234 inc_d();
1235 }
1236 }
1237 inc_advance();
1238 }
1239
1241 template <typename Fragment>
1242 CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment) {
1243 typename PredicateVector::TrivialIterator pred_it;
1244 load_post_increment(fragment, pred_it);
1245 }
1246
1248 template <typename Fragment, typename PredicateIterator>
1249 CUTLASS_HOST_DEVICE void load(Fragment &fragment, PredicateIterator pred_it) const {
1250 TileStoreIterator _load_it(*this);
1251 _load_it.load_post_increment(fragment, pred_it);
1252 }
1253
1255 template <typename Fragment>
1256 CUTLASS_HOST_DEVICE void load(Fragment &fragment) const {
1257 typename PredicateVector::TrivialIterator pred_it;
1258 load(fragment, pred_it);
1259 }
1260
1262 template <typename Fragment>
1263 CUTLASS_HOST_DEVICE void load(Fragment &fragment, int d) {
1264 FragmentIterator frag_iterator(fragment);
1265 for (int h = 0; h < Iterations::kH; ++h) {
1266 for (int w = 0; w < Iterations::kW; ++w) {
1267 for (int c = 0; c < Iterations::kC; ++c) {
1268 load_element(reinterpret_cast<AccessType &>(frag_iterator.at(0, h, w, c)), d, h, w, c);
1269 }
1270 }
1271 }
1272 }
1273 };
1274
1276
1277 } // namespace cutlass
cutlass::TileLoadIterator::inc_advance
CUTLASS_HOST_DEVICE void inc_advance()
Increment in the next dimension.
Definition: tile_iterator.h:689
cutlass::TileStoreIterator::load_element
CUTLASS_HOST_DEVICE void load_element(AccessType &value, int d, int h, int w, int c) const
Loads a single fragment element from memory.
Definition: tile_iterator.h:1198
cutlass::TileIteratorBase::AccessType
Vectorize< FragmentElement, kAccessSize >::Type AccessType
The elements loaded/store by one instruction.
Definition: tile_iterator.h:191
cutlass::TileStoreIterator::Tile
Base::Tile Tile
Tile shape.
Definition: tile_iterator.h:891
cutlass::RegularTilePredicateFunctor::Delta
Delta_ Delta
Definition: tile_iterator.h:113
cutlass::TileStoreIterator::PredicateVector
Base::PredicateVector PredicateVector
Default predicate mask type.
Definition: tile_iterator.h:918
cutlass::TileLoadIterator::Params::Params
CUTLASS_HOST_DEVICE Params()
Initialize params to access storage object.
Definition: tile_iterator.h:507
cutlass::TileLoadIterator::TileLoadIterator
CUTLASS_HOST_DEVICE TileLoadIterator(Params const &_params, Coord< 3 > const &block_offset=make_Coord(0, 0, 0), ThreadOffset thread_offset_func=ThreadOffset())
Constructs a tile load iterator.
Definition: tile_iterator.h:650
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(SharedStorage const &storage)
Initialize params to access storage object.
Definition: tile_iterator.h:550
cutlass::TileTraits::Tile
Tile_ Tile
Shape of the tile.
Definition: tile_iterator.h:80
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: performance_tuning.h:35
cutlass::TileIteratorBase::Index
Index_ Index
Index type.
Definition: tile_iterator.h:164
cutlass
Definition: convert.h:33
tensor_ref.h
Defines a structure containing strides, bounds, and a pointer to tensor data.
cutlass::TileLoadIterator::Iterations
Base::Iterations Iterations
Iterations.
Definition: tile_iterator.h:455
cutlass::TileIteratorBase::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Coord< 4 > const &stride)
Initializes the parameters object from a vector of strides.
Definition: tile_iterator.h:286
cutlass::TileIteratorBase::valid
CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const
Is the iterator valid?
Definition: tile_iterator.h:337
cutlass::TileIteratorBase::Skew
Skew_ Skew
Skew quantity.
Definition: tile_iterator.h:170
cutlass::TileLoadIterator::kAccessSize
static int const kAccessSize
The number of scalars accessed per load/store.
Definition: tile_iterator.h:467
cutlass::MemorySpace
Enum to specify which memory space data resides in.
Definition: load_store.h:38
cutlass::TileStoreIterator::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment)
Loads a fragment and advances the iterator to the next tile.
Definition: tile_iterator.h:1242
cutlass::TileStoreIterator::Index
Base::Index Index
Index type.
Definition: tile_iterator.h:882
cutlass::TileLoadIterator::SharedStorage
Base::Storage SharedStorage
Storage object that may be loaded from.
Definition: tile_iterator.h:482
cutlass::TileStoreIterator::stage
int stage
The stage.
Definition: tile_iterator.h:1028
cutlass::TileLoadIterator::Tile
Base::Tile Tile
Tile shape.
Definition: tile_iterator.h:449
cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::FragmentIterator
FragmentIterator< Fragment, Iterations, AccessType > FragmentIterator
The fragment iterator.
Definition: tile_iterator.h:202
cutlass::TileStoreIterator::Pointer
Scalar * Pointer
Pointer to underlying type.
Definition: tile_iterator.h:927
cutlass::TileIteratorBase::ThreadOffset
Traits::ThreadOffset ThreadOffset
Thread offset.
Definition: tile_iterator.h:185
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
cutlass::IteratorAdvance::Kind
Kind
Definition: tile_iterator.h:65
cutlass::TileStoreIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar *ptr, Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Initializes params.
Definition: tile_iterator.h:998
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:368
cutlass::TileTraits::ImmediateOffsetStrides
Shape< 0, 0, 0, 0 > ImmediateOffsetStrides
Strides for immediate offset computation.
Definition: tile_iterator.h:102
cutlass::TileLoadIterator::FragmentConstIterator
Base::FragmentConstIterator FragmentConstIterator
Fragment const iterator definition.
Definition: tile_iterator.h:476
cutlass::TileTraits
A template defining Tile Traits Concept.
Definition: tile_iterator.h:78
cutlass::TileLoadIterator::FragmentElement
FragmentElement_ FragmentElement
Fragment element.
Definition: tile_iterator.h:428
cutlass::TileLoadIterator::Base
TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > Base
Base class.
Definition: tile_iterator.h:419
cutlass::TileLoadIterator::initialize_predicates
CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it, Coord< 3 > const &bounds, Coord< 3 > const &block_offset=make_Coord(0, 0, 0))
Initializes a predicate vector using a RegularTilePredicateFunctor.
Definition: tile_iterator.h:614
cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::Fragment
Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSize > Fragment
The fragment.
Definition: tile_iterator.h:199
cutlass::TileIteratorBase::Iterations
Traits::Iterations Iterations
Iterations.
Definition: tile_iterator.h:182
cutlass::TileLoadIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment, int d)
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:790
cutlass::TileLoadIterator::Delta
Base::Delta Delta
Delta.
Definition: tile_iterator.h:452
cutlass::TileLoadIterator::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment, PredicateIterator pred_it)
Loads a fragment and advances the iterator to the next tile.
Definition: tile_iterator.h:741
cutlass::FragmentElementType::kWmmaMatrix
Definition: load_store.h:48
cutlass::TileStoreIterator::LongIndex
Base::LongIndex LongIndex
Long index type.
Definition: tile_iterator.h:885
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize()
Definition: tile_iterator.h:590
cutlass::TileLoadIterator::kMemorySpace
static MemorySpace::Kind const kMemorySpace
Source or destination memory space.
Definition: tile_iterator.h:437
cutlass::TileIteratorBase::ImmediateOffsetStrides
Traits::ImmediateOffsetStrides ImmediateOffsetStrides
The strides in each dimension between different loads/stores.
Definition: tile_iterator.h:179
cutlass::TileLoadIterator::ThreadOffset
Base::ThreadOffset ThreadOffset
ThreadOffset functor.
Definition: tile_iterator.h:458
cutlass::TileStoreIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize()
Initializes params to default values.
Definition: tile_iterator.h:1014
cutlass::TileLoadIterator::Fragment
Base::Fragment Fragment
Fragment definition.
Definition: tile_iterator.h:470
cutlass::IteratorAdvance::kH
Definition: tile_iterator.h:65
cutlass::TileStoreIterator::SharedStorage
Base::Storage SharedStorage
Storage object which may be stored to.
Definition: tile_iterator.h:921
cutlass::TileStoreIterator::kMemorySpace
static MemorySpace::Kind const kMemorySpace
Source or destination memory space.
Definition: tile_iterator.h:876
cutlass::TileStoreIterator::initialize_predicates
CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it, Coord< 3 > const &bounds, Coord< 3 > const &block_offset=make_Coord(0, 0, 0))
Initializes a predicate vector using a RegularTilePredicateFunctor.
Definition: tile_iterator.h:1038
cutlass::TileIteratorBase::Params::inc_d
Index inc_d
Definition: tile_iterator.h:226
cutlass::TileStoreIterator::inc_d
CUTLASS_HOST_DEVICE void inc_d()
Increment in the D dimension.
Definition: tile_iterator.h:1098
cutlass::TileTraits::ThreadOffset
ThreadOffset_ ThreadOffset
Functor that returns the logical coordinate of each entity&#39;s initial offset in the tile...
Definition: tile_iterator.h:99
cutlass::PredicateVector::TrivialIterator
Iterator that always returns true.
Definition: predicate_vector.h:309
cutlass::TileIteratorBase::Params::Params
CUTLASS_HOST_DEVICE Params(Coord< 4 > const &stride)
Constructs params with a stride vector.
Definition: tile_iterator.h:259
cutlass::TileStoreIterator::Params::pointer
Scalar * pointer
Pointer to memory.
Definition: tile_iterator.h:935
cutlass::TileStoreIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar *ptr, long long stride_d, Index stride_h, Index stride_w)
Definition: tile_iterator.h:957
cutlass::TileStoreIterator::FragmentShape
Base::FragmentShape FragmentShape
Fragment type.
Definition: tile_iterator.h:903
cutlass::MemorySpace::Kind
Kind
Definition: load_store.h:39
cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::PredicateVector
PredicateVector< ShapeCount< Iterations >::kCount > PredicateVector
Default predicate mask type.
Definition: tile_iterator.h:209
cutlass::TileLoadIterator::stride_advance
CUTLASS_HOST_DEVICE Index stride_advance(void)
Definition: tile_iterator.h:731
cutlass::TileStoreIterator::Fragment
Base::Fragment Fragment
Fragment definition.
Definition: tile_iterator.h:909
cutlass::TileStoreIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar *ptr, Index stride_d, Index stride_h, Index stride_w)
Initializes params to access a raw pointer.
Definition: tile_iterator.h:990
cutlass::TileStoreIterator::TensorRef
TensorRef< Scalar, 4 > TensorRef
Tensor reference for the store iterator.
Definition: tile_iterator.h:930
cutlass::TileLoadIterator::Traits
Base::Traits Traits
concept TileTraits
Definition: tile_iterator.h:422
cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >::TensorRef
TensorRef< Scalar const, 4 > TensorRef
Tensor reference for the load iterator.
Definition: tile_iterator.h:494
cutlass::Store
Definition: load_store.h:178
cutlass::TileIteratorBase::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Initializes params.
Definition: tile_iterator.h:265
cutlass::TileStoreIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment) const
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:1256
cutlass::TileLoadIterator::initialize_predicates
CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it, PredicateFunctor const &functor, Coord< 3 > const &block_offset)
Initializes a predicate vector using an arbitrary predicate functor.
Definition: tile_iterator.h:631
cutlass::TileStoreIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar *ptr, Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Definition: tile_iterator.h:963
cutlass::TileStoreIterator::store_element
CUTLASS_HOST_DEVICE void store_element(AccessType const &value, int d, int h, int w, int c)
Stores a single fragment element into memory.
Definition: tile_iterator.h:1135
cutlass::TileIteratorBase::FragmentShape
FragmentIterator::FragmentShape FragmentShape
The shape of the fragment.
Definition: tile_iterator.h:206
cutlass::TileLoadIterator::Pointer
Scalar const * Pointer
The pointer type.
Definition: tile_iterator.h:491
cutlass::TileLoadIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar const *ptr, Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Initialize params to access storage object.
Definition: tile_iterator.h:521
cutlass::TileIteratorBase::kAdvance
static IteratorAdvance::Kind const kAdvance
Specifies dimension in which post-increment accesses advance.
Definition: tile_iterator.h:155
cutlass::TileIteratorBase::Params::stride_h
Index stride_h
Definition: tile_iterator.h:223
predicate_vector.h
Defines container classes and iterators for managing a statically sized vector of boolean predicates...
cutlass::TileStoreIterator::Params
Parameters.
Definition: tile_iterator.h:933
cutlass::TileIteratorBase::kFragmentElementType
static FragmentElementType::Kind const kFragmentElementType
Specifies iterator storage fragment type (Scalar or WmmaMatrix)
Definition: tile_iterator.h:158
cutlass::TileIteratorBase::Traits
Traits_ Traits
concept TileTraits
Definition: tile_iterator.h:146
cutlass::TileStoreIterator::params
Params params
Parameters structure.
Definition: tile_iterator.h:1022
cutlass::TileStoreIterator::FragmentConstIterator
Base::FragmentConstIterator FragmentConstIterator
Fragment const iterator definition.
Definition: tile_iterator.h:915
cutlass::TileLoadIterator
An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
Definition: tile_iterator.h:402
cutlass::TileStoreIterator::BaseParams
Base::Params BaseParams
IteratorBase parameters.
Definition: tile_iterator.h:924
cutlass::TileLoadIterator::inc_stage
CUTLASS_HOST_DEVICE void inc_stage()
Increment the stage.
Definition: tile_iterator.h:705
cutlass::TileIteratorBase::Params::initialize
CUTLASS_HOST_DEVICE int initialize()
Gotta have this.
Definition: tile_iterator.h:324
cutlass::FragmentElementType::Kind
Kind
Definition: load_store.h:48
cutlass::RegularTilePredicateFunctor::RegularTilePredicateFunctor
CUTLASS_HOST_DEVICE RegularTilePredicateFunctor(Coord< 3 > _bounds)
Constructs a predicate functor given the bounds of a tensor.
Definition: tile_iterator.h:120
cutlass::TileStoreIterator::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset)
Adds a raw offset to the pointer.
Definition: tile_iterator.h:1132
cutlass::TileLoadIterator::TileLoadIterator
CUTLASS_HOST_DEVICE TileLoadIterator()
Default constructor.
Definition: tile_iterator.h:646
cutlass::MemorySpace::kGeneric
Definition: load_store.h:40
cutlass::TileLoadIterator::BaseParams
Base::Params BaseParams
IteratorBase parameters.
Definition: tile_iterator.h:485
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar const *ptr, Index stride_d, Index stride_h, Index stride_w)
Initializes params to access a raw pointer.
Definition: tile_iterator.h:566
cutlass::TileLoadIterator::params
Params params
Parameters structure.
Definition: tile_iterator.h:598
cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::FragmentConstIterator
FragmentConstIterator< Fragment, Iterations, AccessType > FragmentConstIterator
The fragment const iterator.
Definition: tile_iterator.h:204
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar const *ptr)
Initialize params to access storage object.
Definition: tile_iterator.h:558
cutlass::TileLoadIterator::operator+=
CUTLASS_HOST_DEVICE TileLoadIterator & operator+=(Coord< 3 > const &offset)
Adds a vector offset to the iterator.
Definition: tile_iterator.h:719
cutlass::TileLoadIterator::load_element
CUTLASS_HOST_DEVICE void load_element(AccessType &value, int d, int h, int w, int c) const
Loads a single fragment element from memory.
Definition: tile_iterator.h:692
cutlass::TileStoreIterator::AccessType
Base::AccessType AccessType
Memory access type.
Definition: tile_iterator.h:906
cutlass::TileLoadIterator::kRequiresLoadFence
Definition: tile_iterator.h:488
cutlass::DumpType
Definition: tile_iterator.h:134
cutlass::TileTraits::Iterations
Iterations_ Iterations
Number of accesses performed.
Definition: tile_iterator.h:86
cutlass::FragmentIterator::FragmentShape
ShapeMul< Iterations, Shape< 1, 1, 1, kElementsPerAccess > >::Shape FragmentShape
The shape of the the fragment.
Definition: fragment.h:183
cutlass::TileLoadIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment, PredicateIterator pred_it) const
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:776
cutlass::TileLoadIterator::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset)
Adds a raw offset to the pointer.
Definition: tile_iterator.h:729
cutlass::TileTraits::kAccessSize
static int const kAccessSize
Access size.
Definition: tile_iterator.h:105
cutlass::TileIteratorBase::Storage
Fragment< Scalar, ShapeCount< Tile >::kCount, kFragmentSize > Storage
The storage.
Definition: tile_iterator.h:197
cutlass::TileStoreIterator::Skew
Base::Skew Skew
Skew quantity.
Definition: tile_iterator.h:888
cutlass::TileTraits::Delta
Delta_ Delta
Number of steps between accesses along each dimension.
Definition: tile_iterator.h:83
load_store.h
Defines abstractions for efficiently loading and storing vectors to memory.
cutlass::TileLoadIterator::Scalar
Base::Scalar Scalar
Scalar element.
Definition: tile_iterator.h:425
cutlass::TileLoadIterator::FragmentShape
Base::FragmentShape FragmentShape
Fragment type.
Definition: tile_iterator.h:461
cutlass::TileStoreIterator::kFragmentElementType
static FragmentElementType::Kind const kFragmentElementType
Specifies type of iterator fragment storage (Salar or WmmaMatrix)
Definition: tile_iterator.h:873
cutlass::TileStoreIterator::TileStoreIterator
CUTLASS_HOST_DEVICE TileStoreIterator(Params const &, Scalar *ptr, ThreadOffset thread_offset_func=ThreadOffset())
Constructs a tile store iterator.
Definition: tile_iterator.h:1087
cutlass::TileStoreIterator::inc_advance
CUTLASS_HOST_DEVICE void inc_advance()
Increment in the next dimension.
Definition: tile_iterator.h:1107
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::TileStoreIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar *ptr)
Definition: tile_iterator.h:947
cutlass::TileStoreIterator::TileStoreIterator
CUTLASS_HOST_DEVICE TileStoreIterator(Params const &_params, Coord< 3 > const &block_offset=make_Coord(0, 0, 0), ThreadOffset thread_offset_func=ThreadOffset())
Constructs a tile store iterator.
Definition: tile_iterator.h:1074
cutlass::TileLoadIterator::LongIndex
Base::LongIndex LongIndex
Index type.
Definition: tile_iterator.h:443
cutlass::TileIteratorBase::Params::inc_h
Index inc_h
Definition: tile_iterator.h:227
cutlass::TileIteratorBase::Params::Params
CUTLASS_HOST_DEVICE Params()
Constructs params.
Definition: tile_iterator.h:238
cutlass::TileIteratorBase::Params::stride_d
Index stride_d
Definition: tile_iterator.h:222
cutlass::Vector
Definition: vector.h:62
cutlass::TileStoreIterator::Base
TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > Base
Base class.
Definition: tile_iterator.h:858
cutlass::Load
Definition: load_store.h:60
cutlass::TileStoreIterator::Traits
Base::Traits Traits
concept TileTraits
Definition: tile_iterator.h:861
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
cutlass::TileLoadIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar const *ptr)
Initialize params to access storage object.
Definition: tile_iterator.h:511
cutlass::TileStoreIterator::FragmentIterator
Base::FragmentIterator FragmentIterator
Fragment iterator definition.
Definition: tile_iterator.h:912
cutlass::IteratorAdvance
Specifies dimension in which post-increment accesses advance.
Definition: tile_iterator.h:64
cutlass::TileLoadIterator::Params::Params
CUTLASS_HOST_DEVICE Params(TensorRef const &ref)
Constructs with a CompactTensorRef<>
Definition: tile_iterator.h:515
cutlass::TileIteratorBase::kMemorySpace
static MemorySpace::Kind const kMemorySpace
Source or destination memory space.
Definition: tile_iterator.h:161
cutlass::TileStoreIterator::TileStoreIterator
CUTLASS_HOST_DEVICE TileStoreIterator()
Default constructor.
Definition: tile_iterator.h:1070
cutlass::Coord< 3 >
cutlass::FragmentElementType::kScalar
Definition: load_store.h:48
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(TensorRef const &ref)
Initializes params to access a raw pointer.
Definition: tile_iterator.h:543
cutlass::TileStoreIterator::Scalar
Base::Scalar Scalar
Scalar element.
Definition: tile_iterator.h:864
vector.h
Defines a 1D vector of elements held in the registers of each thread.
cutlass::TileIteratorBase
Iterator for accessing a stripmined tile in memory.
Definition: tile_iterator.h:144
cutlass::TileLoadIterator::Params::pointer
Scalar const * pointer
Pointer to memory.
Definition: tile_iterator.h:499
cutlass::TileStoreIterator::store_post_increment
CUTLASS_HOST_DEVICE void store_post_increment(Fragment const &fragment, PredicateIterator pred_it)
Stores a fragment and advances to the next tile.
Definition: tile_iterator.h:1149
cutlass::IteratorAdvance::kD
Definition: tile_iterator.h:65
cutlass::TileLoadIterator::kAdvance
static IteratorAdvance::Kind const kAdvance
Specifies in which dimension post-increment accesses advance.
Definition: tile_iterator.h:431
cutlass::TileIteratorBase::FragmentElement
FragmentElement_ FragmentElement
Fragment element.
Definition: tile_iterator.h:152
cutlass::TileStoreIterator::store
CUTLASS_HOST_DEVICE void store(Fragment const &fragment) const
Stores a fragment without advancing the iterator.
Definition: tile_iterator.h:1192
cutlass::TileStoreIterator::thread_offset
Coord< 4 > thread_offset
Offset of an individual lane from the start of the tile.
Definition: tile_iterator.h:1025
cutlass::TileStoreIterator::Iterations
Base::Iterations Iterations
Iterations.
Definition: tile_iterator.h:897
cutlass::TileStoreIterator::kAccessSize
static int const kAccessSize
The number of scalars accessed per load/store.
Definition: tile_iterator.h:879
cutlass::TileLoadIterator::Params::Params
CUTLASS_HOST_DEVICE Params(Scalar const *ptr, Index stride_d, Index stride_h, Index stride_w)
Initialize params to access storage object.
Definition: tile_iterator.h:536
cutlass::RegularTilePredicateFunctor
Functor computing a predicate given the logical position of an access.
Definition: tile_iterator.h:112
cutlass::TileIteratorBase::Tile
Traits::Tile Tile
Tile shape.
Definition: tile_iterator.h:173
cutlass::TileIteratorBase::kAccessSize
static int const kAccessSize
The number of scalars accessed per load/store.
Definition: tile_iterator.h:188
cutlass::TileLoadIterator::Params
Parameters.
Definition: tile_iterator.h:497
cutlass::ComputeOffsetFromStrides::get
static CUTLASS_HOST_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:199
cutlass::TileLoadIterator::PredicateVector
Base::PredicateVector PredicateVector
Default predicate mask type.
Definition: tile_iterator.h:479
cutlass::TileStoreIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment, int d)
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:1263
cutlass::TileIteratorBase::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Index _stride_d, Index _stride_h, Index _stride_w)
Initializes the parameters object from a vector of strides.
Definition: tile_iterator.h:292
cutlass::TileLoadIterator::AccessType
Base::AccessType AccessType
Memory access type.
Definition: tile_iterator.h:464
cutlass::TileLoadIterator::Skew
Base::Skew Skew
Skew quantity.
Definition: tile_iterator.h:446
cutlass::TileStoreIterator::Delta
Base::Delta Delta
Delta.
Definition: tile_iterator.h:894
cutlass::TileStoreIterator::inc_h
CUTLASS_HOST_DEVICE void inc_h()
Increment in the H dimension.
Definition: tile_iterator.h:1101
cutlass::TileStoreIterator::store
CUTLASS_HOST_DEVICE void store(Fragment const &fragment, PredicateIterator pred_it) const
Stores a fragment without advancing the iterator.
Definition: tile_iterator.h:1185
cutlass::TileLoadIterator::inc_w
CUTLASS_HOST_DEVICE void inc_w()
Increment in the W dimension.
Definition: tile_iterator.h:686
cutlass::TileLoadIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment) const
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:783
cutlass::TileStoreIterator::ThreadOffset
Base::ThreadOffset ThreadOffset
ThreadOffset functor.
Definition: tile_iterator.h:900
cutlass::IteratorAdvance::kW
Definition: tile_iterator.h:65
cutlass::TileLoadIterator::kFragmentElementType
static FragmentElementType::Kind const kFragmentElementType
Specifies type of iterator fragment storage (Salar or WmmaMatrix)
Definition: tile_iterator.h:434
cutlass::TileLoadIterator::inc_h
CUTLASS_HOST_DEVICE void inc_h()
Increment in the H dimension.
Definition: tile_iterator.h:683
cutlass::TileIteratorBase::initialize_predicates
static CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it, PredicateFunctor const &predicate_func, Coord< 3 > const &offset)
Initializes a predicate vector.
Definition: tile_iterator.h:345
cutlass::TileIteratorBase::Scalar
Scalar_ Scalar
Scalar element.
Definition: tile_iterator.h:149
cutlass::TileLoadIterator::thread_offset
Coord< 4 > thread_offset
Offset of an individual lane from the start of the tile.
Definition: tile_iterator.h:601
cutlass::RegularTilePredicateFunctor::bounds
Coord< 3 > bounds
Dimensions of the bounding volume.
Definition: tile_iterator.h:116
cutlass::TileIteratorBase::Delta
Traits::Delta Delta
Distance along each dimension.
Definition: tile_iterator.h:176
cutlass::TileIteratorBase::kFragmentSize
static int const kFragmentSize
The size of storage needed per fragment.
Definition: tile_iterator.h:194
cutlass::TileIteratorBase::Params::inc_advance
Index inc_advance
Definition: tile_iterator.h:230
cutlass::TileStoreIterator::store_post_increment
CUTLASS_HOST_DEVICE void store_post_increment(Fragment const &fragment)
Stores a fragment and advances to the next tile.
Definition: tile_iterator.h:1178
cutlass::TileIteratorBase::LongIndex
long long LongIndex
Long index.
Definition: tile_iterator.h:167
cutlass::TileStoreIterator::Params::Params
CUTLASS_HOST_DEVICE Params()
Definition: tile_iterator.h:943
cutlass::TileStoreIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar *ptr)
Initialize params to access storage object.
Definition: tile_iterator.h:983
cutlass::TileStoreIterator::FragmentElement
Base::FragmentElement FragmentElement
Fragment element.
Definition: tile_iterator.h:867
cutlass::TileStoreIterator::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment, PredicateIterator pred_it) const
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:1249
cutlass::TileStoreIterator::Params::Params
CUTLASS_HOST_DEVICE Params(TensorRef const &ref)
Constructs with a CompactTensorRef<>
Definition: tile_iterator.h:951
cutlass::TileLoadIterator::inc_d
CUTLASS_HOST_DEVICE void inc_d()
Increment in the D dimension.
Definition: tile_iterator.h:680
cutlass::TileLoadIterator::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment)
Loads a fragment and advances the iterator to the next tile.
Definition: tile_iterator.h:769
fragment.h
Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
cutlass::TileIteratorBase::Params
Parameters to the iterator.
Definition: tile_iterator.h:216
cutlass::TileStoreIterator::operator+=
CUTLASS_HOST_DEVICE TileStoreIterator & operator+=(Coord< 3 > const &offset)
Adds a vector offset to the iterator.
Definition: tile_iterator.h:1124
cutlass::TileStoreIterator::inc_stage
CUTLASS_HOST_DEVICE void inc_stage()
Increment the stage.
Definition: tile_iterator.h:1110
cutlass::TileLoadIterator::FragmentIterator
Base::FragmentIterator FragmentIterator
Fragment iterator definition.
Definition: tile_iterator.h:473
cutlass::RegularTilePredicateFunctor::operator()
CUTLASS_HOST_DEVICE bool operator()(Coord< 3 > iteration, Coord< 3 > offset) const
Computes the predicate given the logical position of an access.
Definition: tile_iterator.h:124
cutlass::TileLoadIterator::Index
Base::Index Index
Index type.
Definition: tile_iterator.h:440
cutlass::TileStoreIterator::inc_w
CUTLASS_HOST_DEVICE void inc_w()
Increment in the W dimension.
Definition: tile_iterator.h:1104
cutlass::TileStoreIterator::kAdvance
static IteratorAdvance::Kind const kAdvance
Specifies in which dimension post-increment accesses advance.
Definition: tile_iterator.h:870
cutlass::TileLoadIterator::stage
int stage
Stage argument enables wrapping after some number of tiles have been loaded.
Definition: tile_iterator.h:604
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar const *ptr, Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Initializes params.
Definition: tile_iterator.h:574
cutlass::TileIteratorBase::Params::Params
CUTLASS_HOST_DEVICE Params(Index _stride_d, Index _stride_h, Index _stride_w, Index _inc_d, Index _inc_h, Index _inc_w, Index _inc_advance)
Constructs params.
Definition: tile_iterator.h:242
cutlass::TileLoadIterator::TileLoadIterator
CUTLASS_HOST_DEVICE TileLoadIterator(Params const &, Scalar const *ptr, Coord< 3 > const &block_offset=make_Coord(0, 0, 0), ThreadOffset thread_offset_func=ThreadOffset())
Constructs a tile load iterator.
Definition: tile_iterator.h:665
cutlass::TileStoreIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize(SharedStorage &storage)
Initialize params to access storage object.
Definition: tile_iterator.h:976
cutlass::TileIteratorBase::Params::inc_w
Index inc_w
Definition: tile_iterator.h:228
cutlass::TileStoreIterator::initialize_predicates
CUTLASS_HOST_DEVICE void initialize_predicates(PredicateIterator predicate_it, PredicateFunctor const &functor, Coord< 3 > const &block_offset)
Initializes a predicate vector using an arbitrary predicate functor.
Definition: tile_iterator.h:1055
cutlass::TileStoreIterator
An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
Definition: tile_iterator.h:841
cutlass::TileIteratorBase::Params::stride_w
Index stride_w
Definition: tile_iterator.h:224
cutlass::TileStoreIterator::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment, PredicateIterator pred_it)
Loads a fragment and advances the iterator to the next tile.
Definition: tile_iterator.h:1213