| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

default_mma_core_sm70.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 **************************************************************************************************/
32 #pragma once
33
34 #include "cutlass/cutlass.h"
35 #include "cutlass/array.h"
36
37 #include "cutlass/numeric_types.h"
38 #include "cutlass/matrix_shape.h"
39
40
41 #include "cutlass/layout/tensor_op_multiplicand_sm70.h"
42 #include "cutlass/transform/pitch_linear_thread_map.h"
43 #include "cutlass/transform/threadblock/regular_tile_iterator_tensor_op_sm70.h"
44
45 #include "cutlass/gemm/warp/mma_tensor_op_sm70.h"
46 #include "cutlass/gemm/threadblock/default_mma_core.h"
47
49
50 namespace cutlass {
51 namespace gemm {
52 namespace threadblock {
53
55
63 template <
66 typename Shape_,
68 typename WarpShape_,
70 typename ElementA_,
72 typename ElementB_,
74 typename ElementC_,
76 typename LayoutC_,
78 typename Operator_>
79 struct DefaultMmaCore<Shape_, WarpShape_, GemmShape<8, 8, 4>, ElementA_,
80 layout::ColumnMajor, ElementB_, layout::RowMajor,
81 ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_
82 > {
83 using Shape = Shape_;
84 using WarpShape = WarpShape_;
85 using InstructionShape = GemmShape<8, 8, 4>;
86 using ElementA = ElementA_;
87 using LayoutA = layout::ColumnMajor;
88 using ElementB = ElementB_;
89 using LayoutB = layout::RowMajor;
90 using ElementC = ElementC_;
91 using LayoutC = LayoutC_;
92 using OperatorClass = arch::OpClassTensorOp;
93
95 using Operator = Operator_;
96
98 using WarpCount = GemmShape<
99 Shape::kM / WarpShape::kM,
100 Shape::kN / WarpShape::kN,
101 Shape::kK / WarpShape::kK
102 >;
103
104 // Divisility requirements
105 static_assert(
106 !(Shape::kM % WarpShape::kM) &&
107 !(Shape::kN % WarpShape::kN),
108 "Threadblock-scoped GEMM should be divisible by warp-scoped GEMM size."
109 );
110
112 static int const kWarpSize = warp::WarpSize<arch::OpClassTensorOp>::value;
113
115 static int const kThreads = WarpCount::kCount * kWarpSize;
116
118 static int const kAccessSizeInBits = 128;
119
120 //
121 // Shared memory layouts
122 //
123
124 using SmemLayoutA =
125 layout::ColumnMajorVoltaTensorOpMultiplicandCongruous<
126 sizeof_bits<ElementA>::value>;
127
128 // Shared memory layout
129 using SmemLayoutB =
130 layout::RowMajorVoltaTensorOpMultiplicandBCongruous<
131 sizeof_bits<ElementB>::value>;
132
133 //
134 // Iterators to write to shared memory
135 //
136
138 using IteratorThreadMapA = transform::PitchLinearWarpRakedThreadMap<
139 layout::PitchLinearShape<Shape::kM, Shape::kK>,
140 kThreads,
141 layout::PitchLinearShape<8, 4>,
142 kAccessSizeInBits / sizeof_bits<ElementA>::value
143 >;
144
146 using SmemIteratorA = transform::threadblock::RegularTileIterator<
147 MatrixShape<Shape::kM, Shape::kK>,
148 ElementA,
149 SmemLayoutA,
150 1,
151 IteratorThreadMapA
152 >;
153
155 using IteratorThreadMapB = transform::PitchLinearWarpRakedThreadMap<
156 layout::PitchLinearShape<Shape::kN, Shape::kK>,
157 kThreads,
158 layout::PitchLinearShape<8, 4>,
159 kAccessSizeInBits / sizeof_bits<ElementB>::value
160 >;
161
163 using SmemIteratorB = transform::threadblock::RegularTileIterator<
164 MatrixShape<Shape::kK, Shape::kN>,
165 ElementB,
166 SmemLayoutB,
167 0,
168 IteratorThreadMapB
169 >;
170
171 //
172 // Warp-level matrix multiply operator
173 //
174
175 // Define the warp-level tensor op
176 using Policy = cutlass::gemm::warp::MmaTensorOpPolicy<
177 cutlass::arch::Mma<
178 cutlass::gemm::GemmShape<16, 16, 4>,
179 32,
180 ElementA,
181 LayoutA,
182 ElementB,
183 LayoutB,
184 ElementC,
185 cutlass::layout::RowMajor,
186 cutlass::arch::OpMultiplyAdd
187 >,
188 cutlass::MatrixShape<1, 1>
189 >;
190
191 using MmaTensorOp = cutlass::gemm::warp::MmaVoltaTensorOp<
192 WarpShape,
193 ElementA,
194 SmemLayoutA,
195 ElementB,
196 SmemLayoutB,
197 ElementC,
198 LayoutC,
199 Policy
200 >;
201
203 using MmaPolicy = MmaPolicy<
204 MmaTensorOp,
205 MatrixShape<0, 0>,
206 MatrixShape<0, 0>,
207 WarpCount::kK
208 >;
209 };
210
218 template <
221 typename Shape_,
223 typename WarpShape_,
225 typename ElementA_,
227 typename ElementB_,
229 typename ElementC_,
231 typename LayoutC_,
233 typename Operator_>
234 struct DefaultMmaCore<Shape_, WarpShape_, GemmShape<8, 8, 4>, ElementA_,
235 layout::RowMajor, ElementB_, layout::ColumnMajor,
236 ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_
237 > {
238 using Shape = Shape_;
239 using WarpShape = WarpShape_;
240 using InstructionShape = GemmShape<8, 8, 4>;
241 using ElementA = ElementA_;
242 using LayoutA = layout::RowMajor;
243 using ElementB = ElementB_;
244 using LayoutB = layout::ColumnMajor;
245 using ElementC = ElementC_;
246 using LayoutC = LayoutC_;
247 using OperatorClass = arch::OpClassTensorOp;
248
250 using Operator = Operator_;
251
253 using WarpCount = GemmShape<
254 Shape::kM / WarpShape::kM,
255 Shape::kN / WarpShape::kN,
256 Shape::kK / WarpShape::kK
257 >;
258
259 // Divisility requirements
260 static_assert(
261 !(Shape::kM % WarpShape::kM) &&
262 !(Shape::kN % WarpShape::kN),
263 "Threadblock-scoped GEMM should be divisible by warp-scoped GEMM size."
264 );
265
267 static int const kWarpSize = warp::WarpSize<arch::OpClassTensorOp>::value;
268
270 static int const kThreads = WarpCount::kCount * kWarpSize;
271
273 static int const kAccessSizeInBits = 128;
274
275 //
276 // Shared memory layouts
277 //
278
279 using SmemLayoutA = layout::RowMajorVoltaTensorOpMultiplicandCrosswise<
280 sizeof_bits<ElementA>::value, Shape::kK>;
281
282 // Shared memory layout
283 using SmemLayoutB = layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise<
284 sizeof_bits<ElementB>::value, Shape::kK>;
285
286 //
287 // Iterators to write to shared memory
288 //
289
291 using IteratorThreadMapA = transform::PitchLinearWarpRakedThreadMap<
292 layout::PitchLinearShape<Shape::kK, Shape::kM>,
293 kThreads,
294 layout::PitchLinearShape<4, 8>,
295 kAccessSizeInBits / sizeof_bits<ElementA>::value
296 >;
297
299 using SmemIteratorA = transform::threadblock::RegularTileIterator<
300 MatrixShape<Shape::kM, Shape::kK>,
301 ElementA,
302 SmemLayoutA,
303 0,
304 IteratorThreadMapA
305 >;
306
308 using IteratorThreadMapB = transform::PitchLinearWarpRakedThreadMap<
309 layout::PitchLinearShape<Shape::kK, Shape::kN>,
310 kThreads,
311 layout::PitchLinearShape<4, 8>,
312 kAccessSizeInBits / sizeof_bits<ElementB>::value
313 >;
314
316 using SmemIteratorB = transform::threadblock::RegularTileIterator<
317 MatrixShape<Shape::kK, Shape::kN>,
318 ElementB,
319 SmemLayoutB,
320 1,
321 IteratorThreadMapB
322 >;
323
324 //
325 // Warp-level matrix multiply operator
326 //
327
328 // Define the warp-level tensor op
329 using Policy = cutlass::gemm::warp::MmaTensorOpPolicy<
330 cutlass::arch::Mma<
331 cutlass::gemm::GemmShape<16, 16, 4>,
332 32,
333 ElementA,
334 LayoutA,
335 ElementB,
336 LayoutB,
337 ElementC,
338 cutlass::layout::RowMajor,
339 cutlass::arch::OpMultiplyAdd
340 >,
341 cutlass::MatrixShape<1, 1>
342 >;
343
344 using MmaTensorOp = cutlass::gemm::warp::MmaVoltaTensorOp<
345 WarpShape,
346 ElementA,
347 SmemLayoutA,
348 ElementB,
349 SmemLayoutB,
350 ElementC,
351 LayoutC,
352 Policy
353 >;
354
356 using MmaPolicy = MmaPolicy<
357 MmaTensorOp,
358 MatrixShape<0, 0>,
359 MatrixShape<0, 0>,
360 WarpCount::kK
361 >;
362 };
363
365
373 template <
376 typename Shape_,
378 typename WarpShape_,
380 typename ElementA_,
382 typename ElementB_,
384 typename ElementC_,
386 typename LayoutC_,
388 typename Operator_>
389 struct DefaultMmaCore<Shape_, WarpShape_, GemmShape<8, 8, 4>, ElementA_,
390 layout::RowMajor, ElementB_, layout::RowMajor, ElementC_,
391 LayoutC_, arch::OpClassTensorOp, 2, Operator_
392 > {
393 using Shape = Shape_;
394 using WarpShape = WarpShape_;
395 using InstructionShape = GemmShape<8, 8, 4>;
396 using ElementA = ElementA_;
397 using LayoutA = layout::RowMajor;
398 using ElementB = ElementB_;
399 using LayoutB = layout::RowMajor;
400 using ElementC = ElementC_;
401 using LayoutC = LayoutC_;
402 using OperatorClass = arch::OpClassTensorOp;
403
405 using Operator = Operator_;
406
408 using WarpCount = GemmShape<
409 Shape::kM / WarpShape::kM,
410 Shape::kN / WarpShape::kN,
411 Shape::kK / WarpShape::kK
412 >;
413
414 // Divisility requirements
415 static_assert(
416 !(Shape::kM % WarpShape::kM) &&
417 !(Shape::kN % WarpShape::kN),
418 "Threadblock-scoped GEMM should be divisible by warp-scoped GEMM size."
419 );
420
422 static int const kWarpSize = warp::WarpSize<arch::OpClassTensorOp>::value;
423
425 static int const kThreads = WarpCount::kCount * kWarpSize;
426
428 static int const kAccessSizeInBits = 128;
429
430 //
431 // Shared memory layouts
432 //
433
434 using SmemLayoutA = layout::RowMajorVoltaTensorOpMultiplicandCrosswise<
435 sizeof_bits<ElementA>::value, Shape::kK>;
436
437 // Shared memory layout
438 using SmemLayoutB = layout::RowMajorVoltaTensorOpMultiplicandBCongruous<
439 sizeof_bits<ElementB>::value>;
440
441 //
442 // Iterators to write to shared memory
443 //
444
446 using IteratorThreadMapA = transform::PitchLinearWarpRakedThreadMap<
447 layout::PitchLinearShape<Shape::kK, Shape::kM>,
448 kThreads,
449 layout::PitchLinearShape<4, 8>,
450 kAccessSizeInBits / sizeof_bits<ElementA>::value
451 >;
452
454 using SmemIteratorA = transform::threadblock::RegularTileIterator<
455 MatrixShape<Shape::kM, Shape::kK>,
456 ElementA,
457 SmemLayoutA,
458 0,
459 IteratorThreadMapA
460 >;
461
463 using IteratorThreadMapB = transform::PitchLinearWarpRakedThreadMap<
464 layout::PitchLinearShape<Shape::kN, Shape::kK>,
465 kThreads,
466 layout::PitchLinearShape<8, 4>,
467 kAccessSizeInBits / sizeof_bits<ElementB>::value
468 >;
469
471 using SmemIteratorB = transform::threadblock::RegularTileIterator<
472 MatrixShape<Shape::kK, Shape::kN>,
473 ElementB,
474 SmemLayoutB,
475 0,
476 IteratorThreadMapB
477 >;
478
479 //
480 // Warp-level matrix multiply operator
481 //
482
483 // Define the warp-level tensor op
484 using Policy = cutlass::gemm::warp::MmaTensorOpPolicy<
485 cutlass::arch::Mma<
486 cutlass::gemm::GemmShape<16, 16, 4>,
487 32,
488 ElementA,
489 LayoutA,
490 ElementB,
491 LayoutB,
492 ElementC,
493 cutlass::layout::RowMajor,
494 cutlass::arch::OpMultiplyAdd
495 >,
496 cutlass::MatrixShape<1, 1>
497 >;
498
499 using MmaTensorOp = cutlass::gemm::warp::MmaVoltaTensorOp<
500 WarpShape,
501 ElementA,
502 SmemLayoutA,
503 ElementB,
504 SmemLayoutB,
505 ElementC,
506 LayoutC,
507 Policy
508 >;
509
511 using MmaPolicy = MmaPolicy<
512 MmaTensorOp,
513 MatrixShape<0, 0>,
514 MatrixShape<0, 0>,
515 WarpCount::kK
516 >;
517 };
518
520
528 template <
531 typename Shape_,
533 typename WarpShape_,
535 typename ElementA_,
537 typename ElementB_,
539 typename ElementC_,
541 typename LayoutC_,
543 typename Operator_>
544 struct DefaultMmaCore<Shape_, WarpShape_, GemmShape<8, 8, 4>, ElementA_,
545 layout::ColumnMajor, ElementB_, layout::ColumnMajor,
546 ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_
547 > {
548 using Shape = Shape_;
549 using WarpShape = WarpShape_;
550 using InstructionShape = GemmShape<8, 8, 4>;
551 using ElementA = ElementA_;
552 using LayoutA = layout::ColumnMajor;
553 using ElementB = ElementB_;
554 using LayoutB = layout::ColumnMajor;
555 using ElementC = ElementC_;
556 using LayoutC = LayoutC_;
557 using OperatorClass = arch::OpClassTensorOp;
558
560 using Operator = Operator_;
561
563 using WarpCount = GemmShape<
564 Shape::kM / WarpShape::kM,
565 Shape::kN / WarpShape::kN,
566 Shape::kK / WarpShape::kK
567 >;
568
569 // Divisility requirements
570 static_assert(
571 !(Shape::kM % WarpShape::kM) &&
572 !(Shape::kN % WarpShape::kN),
573 "Threadblock-scoped GEMM should be divisible by warp-scoped GEMM size."
574 );
575
577 static int const kWarpSize = warp::WarpSize<arch::OpClassTensorOp>::value;
578
580 static int const kThreads = WarpCount::kCount * kWarpSize;
581
583 static int const kAccessSizeInBits = 128;
584
585 //
586 // Shared memory layouts
587 //
588
589 using SmemLayoutA = layout::ColumnMajorVoltaTensorOpMultiplicandCongruous<
590 sizeof_bits<ElementA>::value>;
591
592 // Shared memory layout
593 using SmemLayoutB = layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise<
594 sizeof_bits<ElementB>::value, Shape::kK>;
595
596 //
597 // Iterators to write to shared memory
598 //
599
601 using IteratorThreadMapA = transform::PitchLinearWarpRakedThreadMap<
602 layout::PitchLinearShape<Shape::kM, Shape::kK>,
603 kThreads,
604 layout::PitchLinearShape<8, 4>,
605 kAccessSizeInBits / sizeof_bits<ElementA>::value
606 >;
607
609 using SmemIteratorA = transform::threadblock::RegularTileIterator<
610 MatrixShape<Shape::kM, Shape::kK>,
611 ElementA,
612 SmemLayoutA,
613 1,
614 IteratorThreadMapA
615 >;
616
618 using IteratorThreadMapB = transform::PitchLinearWarpRakedThreadMap<
619 layout::PitchLinearShape<Shape::kK, Shape::kN>,
620 kThreads,
621 layout::PitchLinearShape<4, 8>,
622 kAccessSizeInBits / sizeof_bits<ElementB>::value
623 >;
624
626 using SmemIteratorB = transform::threadblock::RegularTileIterator<
627 MatrixShape<Shape::kK, Shape::kN>,
628 ElementB,
629 SmemLayoutB,
630 1,
631 IteratorThreadMapB
632 >;
633
634 //
635 // Warp-level matrix multiply operator
636 //
637
638 // Define the warp-level tensor op
639 using Policy = cutlass::gemm::warp::MmaTensorOpPolicy<
640 cutlass::arch::Mma<
641 cutlass::gemm::GemmShape<16, 16, 4>,
642 32,
643 ElementA,
644 LayoutA,
645 ElementB,
646 LayoutB,
647 ElementC,
648 cutlass::layout::RowMajor,
649 cutlass::arch::OpMultiplyAdd
650 >,
651 cutlass::MatrixShape<1, 1>
652 >;
653
654 using MmaTensorOp = cutlass::gemm::warp::MmaVoltaTensorOp<
655 WarpShape,
656 ElementA,
657 SmemLayoutA,
658 ElementB,
659 SmemLayoutB,
660 ElementC,
661 LayoutC,
662 Policy
663 >;
664
666 using MmaPolicy = MmaPolicy<
667 MmaTensorOp,
668 MatrixShape<0, 0>,
669 MatrixShape<0, 0>,
670 WarpCount::kK
671 >;
672 };
673
674 } // namespace threadblock
675 } // namespace gemm
676 } // namespace cutlass
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::OperatorClass
arch::OpClassTensorOp OperatorClass
Definition: default_mma_core_sm70.h:247
cutlass::layout::RowMajorVoltaTensorOpMultiplicandBCongruous
Template mapping a row-major view of pitch-linear memory to VoltaTensorOpMultiplicandCongruous.
Definition: tensor_op_multiplicand_sm70.h:630
cutlass::MatrixShape
Describes the size of a matrix tile.
Definition: matrix_shape.h:42
cutlass
Definition: aligned_buffer.h:35
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementA
ElementA_ ElementA
Definition: default_mma_core_sm70.h:551
regular_tile_iterator_tensor_op_sm70.h
Templates implementing loading of tiles from pitch-linear rank=2 tensors.
cutlass::gemm::warp::WarpSize
Query the number of threads per warp.
Definition: gemm/warp/mma.h:43
mma_tensor_op_sm70.h
Templates implementing warp-level matrix multiply-accumulate operations targeting Tensor Cores...
cutlass::gemm::threadblock::DefaultMmaCore
Definition: default_mma_core.h:90
pitch_linear_thread_map.h
Templates implementing how threads are mapped to a given tile.
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementA
ElementA_ ElementA
Definition: default_mma_core_sm70.h:396
cutlass::layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise
Definition: tensor_op_multiplicand_sm70.h:848
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementB
ElementB_ ElementB
Definition: default_mma_core_sm70.h:398
cutlass::layout::ColumnMajor
Mapping function for column-major matrices.
Definition: layout/matrix.h:142
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Operator
Operator_ Operator
Default Operator.
Definition: default_mma_core_sm70.h:560
cutlass::layout::PitchLinearShape
Template defining a shape used by pitch-linear operators.
Definition: pitch_linear.h:43
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::OperatorClass
arch::OpClassTensorOp OperatorClass
Definition: default_mma_core_sm70.h:557
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementB
ElementB_ ElementB
Definition: default_mma_core_sm70.h:243
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::OperatorClass
arch::OpClassTensorOp OperatorClass
Definition: default_mma_core_sm70.h:402
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Shape
Shape_ Shape
Definition: default_mma_core_sm70.h:548
matrix_shape.h
Defines a Shape template for matrix tiles.
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::MmaPolicy
MmaPolicy< MmaTensorOp, MatrixShape< 0, 0 >, MatrixShape< 0, 0 >, WarpCount::kK > MmaPolicy
Policy used to define MmaPipelined.
Definition: default_mma_core_sm70.h:671
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::LayoutC
LayoutC_ LayoutC
Definition: default_mma_core_sm70.h:91
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementC
ElementC_ ElementC
Definition: default_mma_core_sm70.h:400
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::MmaPolicy
MmaPolicy< MmaTensorOp, MatrixShape< 0, 0 >, MatrixShape< 0, 0 >, WarpCount::kK > MmaPolicy
Policy used to define MmaPipelined.
Definition: default_mma_core_sm70.h:208
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementB
ElementB_ ElementB
Definition: default_mma_core_sm70.h:88
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Shape
Shape_ Shape
Definition: default_mma_core_sm70.h:83
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementB
ElementB_ ElementB
Definition: default_mma_core_sm70.h:553
default_mma_core.h
Defines basic properties needed by CTA-level GEMMs assuming expectations about data layout of the glo...
cutlass::transform::PitchLinearWarpRakedThreadMap
Definition: pitch_linear_thread_map.h:205
cutlass::gemm::warp::MmaVoltaTensorOp
Structure to compute the matrix product targeting CUDA cores and SIMT math instructions.
Definition: mma_tensor_op_sm70.h:77
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::MmaPolicy
MmaPolicy< MmaTensorOp, MatrixShape< 0, 0 >, MatrixShape< 0, 0 >, WarpCount::kK > MmaPolicy
Policy used to define MmaPipelined.
Definition: default_mma_core_sm70.h:361
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementA
ElementA_ ElementA
Definition: default_mma_core_sm70.h:86
cutlass::gemm::GemmShape
Shape of a matrix multiply-add operation.
Definition: include/cutlass/gemm/gemm.h:57
cutlass::transform::threadblock::RegularTileIterator
Definition: regular_tile_iterator.h:50
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::WarpShape
WarpShape_ WarpShape
Definition: default_mma_core_sm70.h:84
cutlass::gemm::warp::MmaTensorOpPolicy
Policy.
Definition: mma_tensor_op_policy.h:48
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::WarpShape
WarpShape_ WarpShape
Definition: default_mma_core_sm70.h:549
cutlass::layout::RowMajorVoltaTensorOpMultiplicandCrosswise
Definition: tensor_op_multiplicand_sm70.h:943
cutlass::layout::RowMajor
Mapping function for row-major matrices.
Definition: layout/matrix.h:50
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::LayoutC
LayoutC_ LayoutC
Definition: default_mma_core_sm70.h:556
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Operator
Operator_ Operator
Default Operator.
Definition: default_mma_core_sm70.h:250
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::WarpShape
WarpShape_ WarpShape
Definition: default_mma_core_sm70.h:394
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::OperatorClass
arch::OpClassTensorOp OperatorClass
Definition: default_mma_core_sm70.h:92
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementA
ElementA_ ElementA
Definition: default_mma_core_sm70.h:241
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Operator
Operator_ Operator
Default Operator.
Definition: default_mma_core_sm70.h:405
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::MmaPolicy
MmaPolicy< MmaTensorOp, MatrixShape< 0, 0 >, MatrixShape< 0, 0 >, WarpCount::kK > MmaPolicy
Policy used to define MmaPipelined.
Definition: default_mma_core_sm70.h:516
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Operator
Operator_ Operator
Default Operator.
Definition: default_mma_core_sm70.h:95
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementC
ElementC_ ElementC
Definition: default_mma_core_sm70.h:555
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementC
ElementC_ ElementC
Definition: default_mma_core_sm70.h:90
cutlass::arch::Mma
Matrix multiply-add operation.
Definition: arch/mma.h:92
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::LayoutC
LayoutC_ LayoutC
Definition: default_mma_core_sm70.h:401
cutlass::layout::ColumnMajorVoltaTensorOpMultiplicandCongruous
Template mapping a column-major view of pitch-linear memory to VoltaTensorOpMultiplicandCongruous.
Definition: tensor_op_multiplicand_sm70.h:191
tensor_op_multiplicand_sm70.h
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::WarpShape
WarpShape_ WarpShape
Definition: default_mma_core_sm70.h:239
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Shape
Shape_ Shape
Definition: default_mma_core_sm70.h:238
cutlass.h
Basic include for CUTLASS.
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::Shape
Shape_ Shape
Definition: default_mma_core_sm70.h:393
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::LayoutC
LayoutC_ LayoutC
Definition: default_mma_core_sm70.h:246
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 8, 8, 4 >, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ >::ElementC
ElementC_ ElementC
Definition: default_mma_core_sm70.h:245