| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

layout/matrix.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 **************************************************************************************************/
34 #pragma once
35
36 #include "cutlass/cutlass.h"
37 #include "cutlass/matrix_coord.h"
38 #include "cutlass/matrix_traits.h"
39
40 namespace cutlass {
41 namespace layout {
42
44 //
45 // Defines data layouts of various matrix formats usable by TensorRef and other classes.
46 //
48
50 class RowMajor {
51 public:
53 static int const kRank = 2;
54
56 static int const kStrideRank = 1;
57
59 using Index = int32_t;
60
62 using LongIndex = int64_t;
63
65 using TensorCoord = MatrixCoord;
66
68 using Stride = Coord<kStrideRank, Index>;
69
70 private:
71 //
72 // Data members
73 //
74
76 Stride stride_;
77
78 public:
79 //
80 // Methods
81 //
82
84 CUTLASS_HOST_DEVICE
85 RowMajor(Index ldm = 0): stride_(ldm) { }
86
88 CUTLASS_HOST_DEVICE
89 RowMajor(Stride stride): stride_(stride) { }
90
92 CUTLASS_HOST_DEVICE
93 static RowMajor packed(MatrixCoord const &extent) {
94 return RowMajor(extent.column());
95 }
96
99 CUTLASS_HOST_DEVICE
100 LongIndex operator()(MatrixCoord const &coord) const {
101 return coord.row() * stride_[0] + coord.column();
102 }
103
105 CUTLASS_HOST_DEVICE
106 MatrixCoord inverse(LongIndex offset) const {
107 return MatrixCoord(Index(offset / stride_[0]), Index(offset % stride_[0]));
108 }
109
111 CUTLASS_HOST_DEVICE
112 Stride stride() const {
113 return stride_;
114 }
115
117 CUTLASS_HOST_DEVICE
118 Stride & stride() {
119 return stride_;
120 }
121
123 CUTLASS_HOST_DEVICE
124 Index stride(int idx) const {
125 return stride_[idx];
126 }
127
129 CUTLASS_HOST_DEVICE
130 Index & stride(int idx) {
131 return stride_[idx];
132 }
133
135 CUTLASS_HOST_DEVICE
136 LongIndex capacity(MatrixCoord const &extent) const {
137 return extent.row() * stride_[0];
138 }
139 };
140
142 class ColumnMajor {
143 public:
145 static int const kRank = 2;
146
148 static int const kStrideRank = 1;
149
151 using Index = int32_t;
152
154 using LongIndex = int64_t;
155
157 using TensorCoord = MatrixCoord;
158
160 using Stride = Coord<kStrideRank, Index>;
161
162 private:
163 //
164 // Data members
165 //
166
168 Stride stride_;
169
170 public:
171 //
172 // Methods
173 //
174
176 CUTLASS_HOST_DEVICE
177 ColumnMajor(Index ldm = 0): stride_(ldm) { }
178
180 CUTLASS_HOST_DEVICE
181 ColumnMajor(Stride stride): stride_(stride) { }
182
183
185 CUTLASS_HOST_DEVICE
186 static ColumnMajor packed(MatrixCoord const &extent) {
187 return ColumnMajor(extent.row());
188 }
189
192 CUTLASS_HOST_DEVICE
193 LongIndex operator()(MatrixCoord const &coord) const {
194 return coord.row() + coord.column() * stride_[0];
195 }
196
198 CUTLASS_HOST_DEVICE
199 MatrixCoord inverse(LongIndex offset) const {
200 return MatrixCoord(Index(offset % stride_[0]), Index(offset / stride_[0]));
201 }
202
204 CUTLASS_HOST_DEVICE
205 Stride stride() const {
206 return stride_;
207 }
208
210 CUTLASS_HOST_DEVICE
211 Stride & stride() {
212 return stride_;
213 }
214
216 CUTLASS_HOST_DEVICE
217 Index stride(int idx) const {
218 return stride_[idx];
219 }
220
222 CUTLASS_HOST_DEVICE
223 Index & stride(int idx) {
224 return stride_[idx];
225 }
226
228 CUTLASS_HOST_DEVICE
229 LongIndex capacity(MatrixCoord const &extent) const {
230 return extent.column() * stride_[0];
231 }
232 };
233
236 template <int Interleave>
237 struct RowMajorInterleaved {
238
240 static int const kRank = 2;
241
243 static int const kStrideRank = 1;
244
246 using Index = int32_t;
247
249 using LongIndex = int64_t;
250
252 using TensorCoord = MatrixCoord;
253
255 using Stride = Coord<kStrideRank, Index>;
256
258 static int const kInterleave = Interleave;
259
260 private:
261 //
262 // Data members
263 //
264
266 Stride stride_;
267
268 public:
269 //
270 // Methods
271 //
272
274 CUTLASS_HOST_DEVICE
275 RowMajorInterleaved(Index ldm = 0): stride_(ldm) { }
276
278 CUTLASS_HOST_DEVICE
279 RowMajorInterleaved(Stride stride): stride_(stride) { }
280
282 CUTLASS_HOST_DEVICE
283 static RowMajorInterleaved packed(MatrixCoord const &extent) {
284 return RowMajorInterleaved(extent.column() * kInterleave);
285 }
286
289 CUTLASS_HOST_DEVICE
290 LongIndex operator()(MatrixCoord const &coord) const {
291 Index row_major = coord.row() / kInterleave;
292 Index row_minor = coord.row() % kInterleave;
293 return row_major * stride_[0] + coord.column() * kInterleave + row_minor;
294 }
295
297 CUTLASS_HOST_DEVICE
298 MatrixCoord inverse(LongIndex offset) const {
299
300 Index row_major = Index(offset / stride_[0]);
301 Index residual = Index(offset % stride_[0]);
302
303 Index column = residual / kInterleave;
304 Index row_minor = residual % kInterleave;
305
306 return MatrixCoord(row_major * kInterleave + row_minor, column);
307 }
308
310 CUTLASS_HOST_DEVICE
311 Stride stride() const {
312 return stride_;
313 }
314
316 CUTLASS_HOST_DEVICE
317 Stride & stride() {
318 return stride_;
319 }
320
322 CUTLASS_HOST_DEVICE
323 Index stride(int idx) const {
324 return stride_[idx];
325 }
326
328 CUTLASS_HOST_DEVICE
329 Index & stride(int idx) {
330 return stride_[idx];
331 }
332
334 CUTLASS_HOST_DEVICE
335 LongIndex capacity(MatrixCoord const &extent) const {
336 return (extent.row() + kInterleave - 1) / kInterleave * stride_[0];
337 }
338 };
339
342 template <int Interleave>
343 struct ColumnMajorInterleaved {
344
346 static int const kRank = 2;
347
349 static int const kStrideRank = 1;
350
352 using Index = int32_t;
353
355 using LongIndex = int64_t;
356
358 using TensorCoord = MatrixCoord;
359
361 using Stride = Coord<kStrideRank, Index>;
362
364 static int const kInterleave = Interleave;
365
366 private:
367 //
368 // Data members
369 //
370
372 Stride stride_;
373
374 public:
375 //
376 // Methods
377 //
378
380 CUTLASS_HOST_DEVICE
381 ColumnMajorInterleaved(Index ldm = 0): stride_(ldm) { }
382
384 CUTLASS_HOST_DEVICE
385 ColumnMajorInterleaved(Stride stride): stride_(stride) { }
386
387
389 CUTLASS_HOST_DEVICE
390 static ColumnMajorInterleaved packed(MatrixCoord const &extent) {
391 return ColumnMajorInterleaved(extent.row() * kInterleave);
392 }
393
396 CUTLASS_HOST_DEVICE
397 LongIndex operator()(MatrixCoord const &coord) const {
398 Index column_major = coord.column() / kInterleave;
399 Index column_minor = coord.column() % kInterleave;
400 return column_major * stride_[0] + coord.row() * kInterleave + column_minor;
401 }
402
404 CUTLASS_HOST_DEVICE
405 MatrixCoord inverse(LongIndex offset) const {
406
407 Index column_major = Index(offset / stride_[0]);
408 Index residual = Index(offset % stride_[0]);
409
410 Index row = residual / kInterleave;
411 Index column_minor = residual % kInterleave;
412
413 return MatrixCoord(row, column_major * kInterleave + column_minor);
414 }
415
417 CUTLASS_HOST_DEVICE
418 Stride stride() const {
419 return stride_;
420 }
421
423 CUTLASS_HOST_DEVICE
424 Stride & stride() {
425 return stride_;
426 }
427
429 CUTLASS_HOST_DEVICE
430 Index stride(int idx) const {
431 return stride_[idx];
432 }
433
435 CUTLASS_HOST_DEVICE
436 Index & stride(int idx) {
437 return stride_[idx];
438 }
439
441 CUTLASS_HOST_DEVICE
442 LongIndex capacity(MatrixCoord const &extent) const {
443 return (extent.column() + kInterleave - 1) / kInterleave * stride_[0];
444 }
445 };
446
448 enum class Matrix {
449 kColumnMajor,
450 kRowMajor
451 };
452
455 struct ContiguousMatrix {
456
458 static int const kRank = 2;
459
461 static int const kStrideRank = 1;
462
464 using Index = int32_t;
465
467 using LongIndex = int64_t;
468
470 using TensorCoord = MatrixCoord;
471
473 using Stride = Coord<kStrideRank, Index>;
474
475 private:
476 //
477 // Data members
478 //
479
481 Stride stride_;
482
484 Matrix layout_;
485
486 public:
487 //
488 // Methods
489 //
490
492 CUTLASS_HOST_DEVICE
493 ContiguousMatrix(
494 Index ldm = 0,
495 Matrix layout = Matrix::kColumnMajor
496 ):
497 stride_(ldm), layout_(layout) { }
498
500 CUTLASS_HOST_DEVICE
501 static ContiguousMatrix packed(
502 MatrixCoord const &extent,
503 Matrix layout = Matrix::kColumnMajor) {
504
505 Index ldm = 0;
506 if (layout == Matrix::kColumnMajor) {
507 ldm = extent.row();
508 }
509 else if (layout == Matrix::kRowMajor) {
510 ldm = extent.column();
511 }
512 return ContiguousMatrix(ldm, layout);
513 }
514
517 CUTLASS_HOST_DEVICE
518 LongIndex operator()(MatrixCoord const &coord) const {
519 if (layout_ == Matrix::kColumnMajor) {
520 return coord.row() + coord.column() * stride_[0];
521 }
522 else if (layout_ == Matrix::kRowMajor) {
523 return coord.row() * stride_[0] + coord.column();
524 }
525 else {
526 // degenerate case
527 return 0;
528 }
529 }
530
532 CUTLASS_HOST_DEVICE
533 MatrixCoord inverse(LongIndex offset) const {
534 // TODO
535 return MatrixCoord(0, 0);
536 }
537
539 CUTLASS_HOST_DEVICE
540 Stride stride() const {
541 return stride_;
542 }
543
545 CUTLASS_HOST_DEVICE
546 Stride & stride() {
547 return stride_;
548 }
549
551 CUTLASS_HOST_DEVICE
552 Index stride(int idx) const {
553 return stride_[idx];
554 }
555
557 CUTLASS_HOST_DEVICE
558 Index & stride(int idx) {
559 return stride_[idx];
560 }
561
563 CUTLASS_HOST_DEVICE
564 LongIndex capacity(MatrixCoord const &extent) const {
565 if (layout_ == Matrix::kColumnMajor) {
566 return stride_[0] * extent.column();
567 }
568 else if (layout_ == Matrix::kRowMajor) {
569 return stride_[0] * extent.row();
570 }
571 else {
572 // degenerate case
573 return 0;
574 }
575 }
576 };
577
580 template <int BlockRows, int BlockColumns>
581 struct ColumnMajorBlockLinear {
583 static int const kRank = 2;
584
586 static int const kStrideRank = 1;
587
589 using Index = int32_t;
590
592 using LongIndex = int64_t;
593
595 using TensorCoord = MatrixCoord;
596
598 using Stride = Coord<kStrideRank, Index>;
599
601 static int const kBlockRows = BlockRows;
602
604 static int const kBlockColumns = BlockColumns;
605
606 private:
607 //
608 // Data members
609 //
610
612 Stride stride_;
613
614 public:
615 //
616 // Methods
617 //
618
620 CUTLASS_HOST_DEVICE
621 ColumnMajorBlockLinear(Index ldm = 0): stride_(ldm) { }
622
624 CUTLASS_HOST_DEVICE
625 static ColumnMajorBlockLinear packed(MatrixCoord const &extent) {
626 return ColumnMajorBlockLinear(extent.row() * kBlockRows * kBlockColumns);
627 }
628
631 CUTLASS_HOST_DEVICE
632 LongIndex operator()(MatrixCoord const &coord) const {
633 return
634 (coord.row() % kBlockRows) +
635 (coord.column() % kBlockColumns) * kBlockRows +
636 (coord.row() / kBlockRows) * kBlockRows * kBlockColumns +
637 (coord.column() / kBlockColumns) * stride_[0];
638 }
639
641 CUTLASS_HOST_DEVICE
642 MatrixCoord inverse(LongIndex offset) const {
643
644 // TODO
645 return MatrixCoord(0, 0);
646 }
647
649 CUTLASS_HOST_DEVICE
650 Stride stride() const {
651 return stride_;
652 }
653
655 CUTLASS_HOST_DEVICE
656 Stride & stride() {
657 return stride_;
658 }
659
661 CUTLASS_HOST_DEVICE
662 Index stride(int idx) const {
663 return stride_[idx];
664 }
665
667 CUTLASS_HOST_DEVICE
668 Index & stride(int idx) {
669 return stride_[idx];
670 }
671
673 CUTLASS_HOST_DEVICE
674 LongIndex capacity(MatrixCoord const &extent) const {
675 return (extent.column() + kBlockColumns - 1) / kBlockColumns * stride_[0];
676 }
677 };
678
681 template <int BlockRows, int BlockColumns>
682 struct RowMajorBlockLinear {
684 static int const kRank = 2;
685
687 static int const kStrideRank = 1;
688
690 using Index = int32_t;
691
693 using LongIndex = int64_t;
694
696 using TensorCoord = MatrixCoord;
697
699 using Stride = Coord<kStrideRank, Index>;
700
702 static int const kBlockRows = BlockRows;
703
705 static int const kBlockColumns = BlockColumns;
706
707 private:
708 //
709 // Data members
710 //
711
713 Stride stride_;
714
715 public:
716 //
717 // Methods
718 //
719
721 CUTLASS_HOST_DEVICE
722 RowMajorBlockLinear(Index ldm = 0): stride_(ldm) { }
723
725 CUTLASS_HOST_DEVICE
726 static RowMajorBlockLinear packed(MatrixCoord const &extent) {
727 return RowMajorBlockLinear(extent.column() * kBlockRows * kBlockColumns);
728 }
729
732 CUTLASS_HOST_DEVICE
733 LongIndex operator()(MatrixCoord const &coord) const {
734 return
735 (coord.column() % kBlockColumns) +
736 (coord.row() % kBlockRows) * kBlockColumns +
737 (coord.column() / kBlockColumns) * kBlockRows * kBlockColumns +
738 (coord.row() / kBlockRows) * stride_[0];
739 }
740
742 CUTLASS_HOST_DEVICE
743 MatrixCoord inverse(LongIndex offset) const {
744 // TODO
745 return MatrixCoord(0, 0);
746 }
747
749 CUTLASS_HOST_DEVICE
750 Stride stride() const {
751 return stride_;
752 }
753
755 CUTLASS_HOST_DEVICE
756 Stride & stride() {
757 return stride_;
758 }
759
761 CUTLASS_HOST_DEVICE
762 Index stride(int idx) const {
763 return stride_[idx];
764 }
765
767 CUTLASS_HOST_DEVICE
768 Index & stride(int idx) {
769 return stride_[idx];
770 }
771
773 CUTLASS_HOST_DEVICE
774 LongIndex capacity(MatrixCoord const &extent) const {
775 return (extent.row() + kBlockRows - 1) / kBlockRows * stride_[0];
776 }
777 };
778
780
781 struct GeneralMatrix {
782
784 static int const kRank = 2;
785
787 static int const kStrideRank = 2;
788
790 using Index = int32_t;
791
793 using LongIndex = int64_t;
794
796 using TensorCoord = MatrixCoord;
797
799 using Stride = Coord<kStrideRank, Index>;
800
801 private:
802 //
803 // Data members
804 //
805
806 MatrixLayout layout_id_;
807
809 Stride stride_;
810
811 public:
812 //
813 // Methods
814 //
815
817 CUTLASS_HOST_DEVICE
818 GeneralMatrix(): layout_id_(MatrixLayout::kColumnMajor), stride_(make_Coord(0, 1)) { }
819
821 CUTLASS_HOST_DEVICE
822 GeneralMatrix(
823 MatrixLayout layout_id,
824 Index ldm,
825 Index interleave): layout_id_(layout_id), stride_(make_Coord(ldm, interleave)) { }
826
828 CUTLASS_HOST_DEVICE
829 static GeneralMatrix packed(
830 MatrixCoord const &extent,
831 MatrixLayout layout_id = MatrixLayout::kColumnMajor,
832 Index interleave = 1) {
833
834 Index c;
835 if (layout_id == MatrixLayout::kRowMajor) {
836 c = extent.column();
837 }
838 else {
839 c = extent.row();
840 }
841
842 Index ldm = c * interleave;
843
844 return GeneralMatrix(layout_id, ldm, interleave);
845 }
846
849 CUTLASS_HOST_DEVICE
850 LongIndex operator()(MatrixCoord const &coord) const {
851 Index c, s;
852 if (layout_id_ == MatrixLayout::kRowMajor) {
853 c = coord.column();
854 s = coord.row();
855 }
856 else {
857 s = coord.column();
858 c = coord.row();
859 }
860
861 Index v = s / stride_[1];
862 Index residual = (s % stride_[1]);
863
864 return LongIndex(c) * LongIndex(stride_[1]) + LongIndex(v) * LongIndex(stride_[0]) + residual;
865 }
866
868 CUTLASS_HOST_DEVICE
869 Stride stride() const {
870 return stride_;
871 }
872
873 CUTLASS_HOST_DEVICE
874 MatrixLayout layout_id() const {
875 return layout_id_;
876 }
877
879 CUTLASS_HOST_DEVICE
880 Stride & stride() {
881 return stride_;
882 }
883
884 CUTLASS_HOST_DEVICE
885 MatrixLayout & layout_id() {
886 return layout_id_;
887 }
888
890 CUTLASS_HOST_DEVICE
891 Index stride(int idx) const {
892 return stride_[idx];
893 }
894
896 CUTLASS_HOST_DEVICE
897 Index & stride(int idx) {
898 return stride_[idx];
899 }
900
902 CUTLASS_HOST_DEVICE
903 LongIndex capacity(MatrixCoord const &extent) const {
904 Index s;
905 if (layout_id_ == MatrixLayout::kRowMajor) {
906 s = extent.row();
907 }
908 else {
909 s = extent.column();
910 }
911
912 Index v = Index((s + stride_[1] - 1) / stride_[1]);
913 return LongIndex(v) * LongIndex(stride_[0]);
914 }
915 };
916
918
920 template <typename Layout>
921 struct LayoutTranspose;
922
924 template <>
925 struct LayoutTranspose<layout::RowMajor> {
926 using type = layout::ColumnMajor;
927 };
928
930 template <>
931 struct LayoutTranspose<layout::ColumnMajor> {
932 using type = layout::RowMajor;
933 };
934
936
937 } // namespace layout
938 } // namespace cutlass
cutlass::layout::ContiguousMatrix::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:464
cutlass::layout::RowMajorInterleaved::packed
static CUTLASS_HOST_DEVICE RowMajorInterleaved packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:283
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:323
cutlass::layout::RowMajor::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:62
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:668
cutlass::layout::ColumnMajorInterleaved< 4 >::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:355
cutlass::MatrixCoord::column
CUTLASS_HOST_DEVICE Index const & column() const
Returns the column of the coordinate.
Definition: matrix_coord.h:85
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:540
cutlass::layout::RowMajorBlockLinear::packed
static CUTLASS_HOST_DEVICE RowMajorBlockLinear packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:726
cutlass
Definition: aligned_buffer.h:35
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:750
cutlass::layout::Matrix::kColumnMajor
leading dimension refers to stride between columns; stride along rows is 1
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:311
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:768
cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved
CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Index ldm=0)
Ctor.
Definition: layout/matrix.h:381
cutlass::layout::GeneralMatrix::GeneralMatrix
CUTLASS_HOST_DEVICE GeneralMatrix(MatrixLayout layout_id, Index ldm, Index interleave)
Ctor.
Definition: layout/matrix.h:822
cutlass::layout::ColumnMajorBlockLinear::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:592
cutlass::layout::ContiguousMatrix
Definition: layout/matrix.h:455
cutlass::layout::GeneralMatrix
Definition: layout/matrix.h:781
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:662
cutlass::layout::ColumnMajorInterleaved::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:442
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:130
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:124
cutlass::layout::ContiguousMatrix::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:467
cutlass::layout::RowMajorInterleaved< 4 >::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:249
cutlass::layout::ColumnMajorInterleaved::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:397
cutlass::layout::GeneralMatrix::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:793
cutlass::layout::RowMajorBlockLinear::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:693
cutlass::layout::RowMajor::kRank
static int const kRank
Logical rank of tensor.
Definition: layout/matrix.h:53
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:211
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:329
cutlass::layout::ColumnMajorInterleaved< 4 >::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:352
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:558
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:387
cutlass::layout::ColumnMajor::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:154
cutlass::layout::ColumnMajor::ColumnMajor
CUTLASS_HOST_DEVICE ColumnMajor(Index ldm=0)
Ctor.
Definition: layout/matrix.h:177
cutlass::layout::RowMajorInterleaved::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:298
cutlass::layout::RowMajorBlockLinear::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:733
cutlass::layout::RowMajorInterleaved< 4 >::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:246
cutlass::MatrixCoord::row
CUTLASS_HOST_DEVICE Index const & row() const
Returns the row of the coordinate.
Definition: matrix_coord.h:77
cutlass::layout::RowMajor::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:136
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:418
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:112
cutlass::layout::RowMajorInterleaved::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:335
cutlass::layout::RowMajor::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:106
cutlass::layout::ColumnMajorBlockLinear
Definition: layout/matrix.h:581
cutlass::layout::ColumnMajor
Mapping function for column-major matrices.
Definition: layout/matrix.h:142
cutlass::layout::GeneralMatrix::layout_id
CUTLASS_HOST_DEVICE MatrixLayout & layout_id()
Definition: layout/matrix.h:885
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:762
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:891
cutlass::layout::GeneralMatrix::GeneralMatrix
CUTLASS_HOST_DEVICE GeneralMatrix()
Ctor.
Definition: layout/matrix.h:818
cutlass::layout::RowMajor::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:59
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:650
cutlass::layout::ColumnMajorBlockLinear::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:642
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:897
cutlass::layout::RowMajor::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:100
cutlass::layout::GeneralMatrix::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:790
cutlass::layout::LayoutTranspose
Defines transposes of matrix layouts.
Definition: layout/matrix.h:921
cutlass::layout::ColumnMajorBlockLinear::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:632
cutlass::layout::ContiguousMatrix::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:564
cutlass::layout::ColumnMajorBlockLinear::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:589
cutlass::layout::Matrix
Matrix
Enumerated type for canonical pitch-linear matrix layouts.
Definition: layout/matrix.h:448
cutlass::layout::ContiguousMatrix::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:518
cutlass::layout::ColumnMajor::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:199
cutlass::layout::RowMajorBlockLinear
Definition: layout/matrix.h:682
cutlass::MatrixLayout
MatrixLayout
Definition: matrix_traits.h:36
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:436
cutlass::layout::RowMajorBlockLinear::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:690
cutlass::layout::RowMajorInterleaved::RowMajorInterleaved
CUTLASS_HOST_DEVICE RowMajorInterleaved(Stride stride)
Ctor.
Definition: layout/matrix.h:279
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::layout::ContiguousMatrix::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:533
cutlass::layout::RowMajorBlockLinear::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:743
cutlass::layout::ColumnMajor::packed
static CUTLASS_HOST_DEVICE ColumnMajor packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:186
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:756
cutlass::layout::RowMajorBlockLinear::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:774
cutlass::layout::GeneralMatrix::packed
static CUTLASS_HOST_DEVICE GeneralMatrix packed(MatrixCoord const &extent, MatrixLayout layout_id=MatrixLayout::kColumnMajor, Index interleave=1)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:829
cutlass::layout::ColumnMajor::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:229
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:118
cutlass::layout::ContiguousMatrix::packed
static CUTLASS_HOST_DEVICE ContiguousMatrix packed(MatrixCoord const &extent, Matrix layout=Matrix::kColumnMajor)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:501
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:205
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:424
cutlass::layout::Matrix::kRowMajor
leading dimension refers to stride between rows; stride along columns is 1
cutlass::Coord< kStrideRank, Index >
cutlass::layout::RowMajor
Mapping function for row-major matrices.
Definition: layout/matrix.h:50
cutlass::layout::ColumnMajor::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:193
matrix_coord.h
Defines a canonical coordinate for rank=2 matrices offering named indices.
cutlass::layout::RowMajor::RowMajor
CUTLASS_HOST_DEVICE RowMajor(Stride stride)
Ctor.
Definition: layout/matrix.h:89
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:880
cutlass::layout::RowMajorInterleaved::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:290
cutlass::layout::ContiguousMatrix::ContiguousMatrix
CUTLASS_HOST_DEVICE ContiguousMatrix(Index ldm=0, Matrix layout=Matrix::kColumnMajor)
Ctor.
Definition: layout/matrix.h:493
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:430
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:317
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:552
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:217
cutlass::layout::ColumnMajorBlockLinear::packed
static CUTLASS_HOST_DEVICE ColumnMajorBlockLinear packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:625
cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved
CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Stride stride)
Ctor.
Definition: layout/matrix.h:385
cutlass::layout::GeneralMatrix::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:850
cutlass::layout::ColumnMajorInterleaved
Definition: layout/matrix.h:343
cutlass::layout::ColumnMajor::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:151
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:869
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:656
cutlass::layout::RowMajor::packed
static CUTLASS_HOST_DEVICE RowMajor packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:93
cutlass::layout::RowMajorInterleaved::RowMajorInterleaved
CUTLASS_HOST_DEVICE RowMajorInterleaved(Index ldm=0)
Ctor.
Definition: layout/matrix.h:275
cutlass::layout::RowMajorBlockLinear::RowMajorBlockLinear
CUTLASS_HOST_DEVICE RowMajorBlockLinear(Index ldm=0)
Ctor.
Definition: layout/matrix.h:722
cutlass::layout::GeneralMatrix::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:903
cutlass::layout::ColumnMajorInterleaved::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:405
matrix_traits.h
Defines properties of matrices used to denote layout and operands to GEMM kernels.
cutlass::layout::ColumnMajorInterleaved::packed
static CUTLASS_HOST_DEVICE ColumnMajorInterleaved packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:390
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:223
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:546
cutlass.h
Basic include for CUTLASS.
cutlass::MatrixCoord
Definition: matrix_coord.h:39
cutlass::layout::ColumnMajorBlockLinear::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:674
cutlass::layout::ColumnMajor::ColumnMajor
CUTLASS_HOST_DEVICE ColumnMajor(Stride stride)
Ctor.
Definition: layout/matrix.h:181
cutlass::layout::RowMajor::kStrideRank
static int const kStrideRank
Rank of stride vector.
Definition: layout/matrix.h:56
cutlass::layout::ColumnMajorBlockLinear::ColumnMajorBlockLinear
CUTLASS_HOST_DEVICE ColumnMajorBlockLinear(Index ldm=0)
Ctor.
Definition: layout/matrix.h:621
cutlass::layout::RowMajor::RowMajor
CUTLASS_HOST_DEVICE RowMajor(Index ldm=0)
Constructor.
Definition: layout/matrix.h:85
cutlass::layout::GeneralMatrix::layout_id
CUTLASS_HOST_DEVICE MatrixLayout layout_id() const
Definition: layout/matrix.h:874
cutlass::layout::RowMajorInterleaved
Definition: layout/matrix.h:237