| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

default_gemm.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 **************************************************************************************************/
25
36 #pragma once
37
38 #include "cutlass/cutlass.h"
39
40 #include "cutlass/layout/matrix.h"
41 #include "cutlass/numeric_types.h"
42 #include "cutlass/arch/wmma.h"
43
44 #include "cutlass/epilogue/threadblock/epilogue.h"
45 #include "cutlass/epilogue/thread/linear_combination.h"
46
47 #include "cutlass/gemm/gemm.h"
48 #include "cutlass/gemm/kernel/gemm.h"
49 #include "cutlass/gemm/kernel/gemm_pipelined.h"
50 #include "cutlass/gemm/threadblock/default_mma_core_sm75.h"
51 #include "cutlass/gemm/threadblock/default_mma_core_sm70.h"
52 #include "cutlass/gemm/threadblock/default_mma.h"
53 #include "cutlass/gemm/threadblock/default_mma_core_simt.h"
54 #include "cutlass/gemm/threadblock/threadblock_swizzle.h"
55
56 #include "cutlass/epilogue/threadblock/default_epilogue_tensor_op.h"
57 #include "cutlass/epilogue/threadblock/default_epilogue_volta_tensor_op.h"
58 #include "cutlass/epilogue/threadblock/default_epilogue_simt.h"
59 #include "cutlass/transform/threadblock/predicated_tile_iterator.h"
60
61 #if defined(CUTLASS_ARCH_WMMA_ENABLED)
62 #include "cutlass/epilogue/threadblock/default_epilogue_wmma_tensor_op.h"
63 #endif //CUTLASS_ARCH_WMMA_ENABLED
64
65
67
68 namespace cutlass {
69 namespace gemm {
70 namespace kernel {
71
73
74 template <
76 typename ElementA_,
78 typename LayoutA_,
80 int kAlignmentA,
82 typename ElementB_,
84 typename LayoutB_,
86 int kAlignmentB,
88 typename ElementC_,
90 typename LayoutC_,
92 typename ElementAccumulator,
94 typename OperatorClass,
96 typename ArchTag,
98 typename ThreadblockShape,
100 typename WarpShape,
102 typename InstructionShape,
104 typename EpilogueOutputOp,
106 typename ThreadblockSwizzle,
108 int Stages,
111 bool SplitKSerial,
113 typename Operator,
115 bool IsBetaZero = false>
116 struct DefaultGemm;
117
120 template <
122 typename ElementA,
124 typename LayoutA,
126 int kAlignmentA,
128 typename ElementB,
130 typename LayoutB,
132 int kAlignmentB,
134 typename ElementC,
136 typename ElementAccumulator,
138 typename ThreadblockShape,
140 typename WarpShape,
142 typename InstructionShape,
144 typename EpilogueOutputOp,
146 typename ThreadblockSwizzle,
148 bool SplitKSerial,
150 typename Operator
151 >
152 struct DefaultGemm<
153 ElementA, LayoutA, kAlignmentA,
154 ElementB, LayoutB, kAlignmentB,
155 ElementC, layout::RowMajor,
156 ElementAccumulator,
157 arch::OpClassTensorOp,
158 arch::Sm75,
159 ThreadblockShape,
160 WarpShape,
161 InstructionShape,
162 EpilogueOutputOp,
163 ThreadblockSwizzle,
164 2,
165 SplitKSerial,
166 Operator
167 > {
168
170 using Mma = typename cutlass::gemm::threadblock::DefaultMma<
171 ElementA,
172 LayoutA,
173 kAlignmentA,
174 ElementB,
175 LayoutB,
176 kAlignmentB,
177 ElementAccumulator,
178 layout::RowMajor,
179 arch::OpClassTensorOp,
180 arch::Sm75,
181 ThreadblockShape,
182 WarpShape,
183 InstructionShape,
184 2,
185 Operator
186 >::ThreadblockMma;
187
188 static const int kPartitionsK = ThreadblockShape::kK / WarpShape::kK;
189
191 using Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueTensorOp<
192 ThreadblockShape,
193 typename Mma::Operator,
194 kPartitionsK,
195 EpilogueOutputOp,
196 EpilogueOutputOp::kCount
197 >::Epilogue;
198
200 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
201 };
202
205 template <
207 typename ElementA,
209 int kAlignmentA,
211 typename ElementB,
213 int kAlignmentB,
215 typename ElementC,
217 typename ThreadblockShape,
219 typename WarpShape,
221 typename InstructionShape,
223 typename EpilogueOutputOp,
225 typename ThreadblockSwizzle,
227 int InterleavedK,
230 bool SplitKSerial,
232 typename Operator,
234 bool IsBetaZero>
235 struct DefaultGemm<ElementA, layout::ColumnMajorInterleaved<InterleavedK>,
236 kAlignmentA, ElementB,
237 layout::RowMajorInterleaved<InterleavedK>, kAlignmentB,
238 ElementC, layout::ColumnMajorInterleaved<InterleavedK>,
239 int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape,
240 WarpShape, InstructionShape, EpilogueOutputOp,
241 ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero> {
242 using LayoutA = layout::ColumnMajorInterleaved<InterleavedK>;
243 using LayoutB = layout::RowMajorInterleaved<InterleavedK>;
244 using LayoutC = layout::ColumnMajorInterleaved<InterleavedK>;
245
246 using ElementAccumulator = int32_t;
247
249 using Mma = typename cutlass::gemm::threadblock::DefaultMma<
250 ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC,
251 arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape,
252 InstructionShape, 2, Operator, true>::ThreadblockMma;
253
254 static const int kPartitionsK = ThreadblockShape::kK / WarpShape::kK;
255
257 using Epilogue = typename cutlass::epilogue::threadblock::
258 DefaultInterleavedEpilogueTensorOp<
259 ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp,
260 64 / sizeof_bits<ElementC>::value, InterleavedK,
261 IsBetaZero>::Epilogue;
262
264 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
265 };
266
268
269
271 template <
273 typename ElementA,
275 typename LayoutA,
277 int kAlignmentA,
279 typename ElementB,
281 typename LayoutB,
283 int kAlignmentB,
285 typename ElementC,
287 typename ElementAccumulator,
289 typename ThreadblockShape,
291 typename WarpShape,
293 typename EpilogueOutputOp,
295 typename ThreadblockSwizzle,
297 bool SplitKSerial,
299 typename Operator
300 >
301 struct DefaultGemm<
302 ElementA, LayoutA, kAlignmentA,
303 ElementB, LayoutB, kAlignmentB,
304 ElementC, layout::RowMajor,
305 ElementAccumulator,
306 arch::OpClassTensorOp,
307 arch::Sm70,
308 ThreadblockShape,
309 WarpShape,
310 GemmShape<8, 8, 4>,
311 EpilogueOutputOp,
312 ThreadblockSwizzle,
313 2,
314 SplitKSerial,
315 Operator
316 > {
317
319 using Mma = typename cutlass::gemm::threadblock::DefaultMma<
320 ElementA,
321 LayoutA,
322 kAlignmentA,
323 ElementB,
324 LayoutB,
325 kAlignmentB,
326 ElementAccumulator,
327 layout::RowMajor,
328 arch::OpClassTensorOp,
329 arch::Sm70,
330 ThreadblockShape,
331 WarpShape,
332 GemmShape<8, 8, 4>,
333 2,
334 Operator
335 >::ThreadblockMma;
336
337 static const int kPartitionsK = ThreadblockShape::kK / WarpShape::kK;
338
340 using Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp<
341 ThreadblockShape,
342 typename Mma::Operator,
343 kPartitionsK,
344 EpilogueOutputOp,
345 EpilogueOutputOp::kCount
346 >::Epilogue;
347
349 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
350 };
351
353
355 template <
357 typename ElementA,
359 typename LayoutA,
361 int kAlignmentA,
363 typename ElementB,
365 typename LayoutB,
367 int kAlignmentB,
369 typename ElementC,
371 typename ElementAccumulator,
373 typename ArchTag,
375 typename ThreadblockShape,
377 typename WarpShape,
379 typename EpilogueOutputOp,
381 typename ThreadblockSwizzle,
383 bool SplitKSerial,
385 typename Operator
386 >
387 struct DefaultGemm<
388 ElementA,
389 LayoutA,
390 kAlignmentA,
391 ElementB,
392 LayoutB,
393 kAlignmentB,
394 ElementC,
395 layout::RowMajor,
396 ElementAccumulator,
397 arch::OpClassSimt,
398 ArchTag,
399 ThreadblockShape,
400 WarpShape,
401 GemmShape<1, 1, 1>,
402 EpilogueOutputOp,
403 ThreadblockSwizzle,
404 2,
405 SplitKSerial,
406 Operator> {
408 using Mma = typename cutlass::gemm::threadblock::DefaultMma<
409 ElementA,
410 LayoutA,
411 kAlignmentA,
412 ElementB,
413 LayoutB,
414 kAlignmentB,
415 ElementAccumulator,
416 layout::RowMajor,
417 arch::OpClassSimt,
418 arch::Sm50,
419 ThreadblockShape,
420 WarpShape,
421 GemmShape<1, 1, 1>,
422 2,
423 Operator>::ThreadblockMma;
424
425 static int const kEpilogueElementsPerAccess = EpilogueOutputOp::kCount;
426 static_assert(kEpilogueElementsPerAccess == 1, "simt epilogue must operate on scalars");
427
429 using Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueSimt<
430 ThreadblockShape,
431 typename Mma::Operator,
432 EpilogueOutputOp,
433 kEpilogueElementsPerAccess
434 >::Epilogue;
435
437 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
438 };
439
441
444
445 template <
447 typename LayoutA,
449 int kAlignmentA,
451 typename LayoutB,
453 int kAlignmentB,
455 typename LayoutC,
457 typename ElementC,
459 typename ArchTag,
461 typename ElementAccumulator,
463 typename ThreadblockShape,
465 typename WarpShape,
467 typename EpilogueOutputOp,
469 typename ThreadblockSwizzle,
472 bool SplitKSerial,
474 typename Operator>
475 struct DefaultGemm<int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB,
476 ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt,
477 ArchTag, ThreadblockShape, WarpShape, GemmShape<1, 1, 4>,
478 EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial,
479 Operator, false> {
480 using InstructionShape = GemmShape<1, 1, 4>;
481 using ElementA = int8_t;
482 using ElementB = int8_t;
483
484 using OperatorClass = arch::OpClassSimt;
486 using Mma = typename cutlass::gemm::threadblock::DefaultMma<ElementA,
487 LayoutA,
488 kAlignmentA,
489 ElementB,
490 LayoutB,
491 kAlignmentB,
492 ElementAccumulator,
493 LayoutC,
494 arch::OpClassSimt,
495 arch::Sm50,
496 ThreadblockShape,
497 WarpShape,
498 InstructionShape,
499 2,
500 Operator,
501 false
502 >::ThreadblockMma;
503
504 static int const kEpilogueElementsPerAccess = EpilogueOutputOp::kCount;
505 static_assert(kEpilogueElementsPerAccess == 1, "simt epilogue must operate on scalars");
506
508 using Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueSimt<
509 ThreadblockShape,
510 typename Mma::Operator,
511 EpilogueOutputOp,
512 kEpilogueElementsPerAccess
513 >::Epilogue;
514
516 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
517 };
518
519
520 #if defined(CUTLASS_ARCH_WMMA_ENABLED)
521 template <
525 typename ElementA,
527 typename LayoutA,
529 int kAlignmentA,
531 typename ElementB,
533 typename LayoutB,
535 int kAlignmentB,
537 typename ElementC,
539 typename LayoutC,
541 typename ElementAccumulator,
543 typename ArchTag,
545 typename ThreadblockShape,
547 typename WarpShape,
549 typename InstructionShape,
551 typename EpilogueOutputOp,
553 typename ThreadblockSwizzle,
555 int Stages,
558 bool SplitKSerial,
560 typename Operator>
561 struct DefaultGemm<
562 ElementA, LayoutA, kAlignmentA,
563 ElementB, LayoutB, kAlignmentB,
564 ElementC, LayoutC,
565 ElementAccumulator,
566 arch::OpClassWmmaTensorOp,
567 ArchTag,
568 ThreadblockShape, WarpShape, InstructionShape,
569 EpilogueOutputOp,
570 ThreadblockSwizzle,
571 Stages,
572 SplitKSerial,
573 Operator> {
575 using Mma = typename cutlass::gemm::threadblock::DefaultMma<
576 ElementA, LayoutA, kAlignmentA,
577 ElementB, LayoutB, kAlignmentB,
578 ElementAccumulator, LayoutC,
579 arch::OpClassWmmaTensorOp,
580 ArchTag,
581 ThreadblockShape,
582 WarpShape,
583 InstructionShape,
584 Stages,
585 Operator>::ThreadblockMma;
586
587 static const int kPartitionsK = ThreadblockShape::kK / WarpShape::kK;
588
590 using Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueWmmaTensorOp<
591 ThreadblockShape,
592 typename Mma::Operator,
593 kPartitionsK,
594 EpilogueOutputOp,
595 EpilogueOutputOp::kCount
596 >::Epilogue;
597
599 using GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>;
600 };
602 #endif //CUTLASS_ARCH_WMMA_ENABLED
603
605
606 } // namespace kernel
607 } // namespace gemm
608 } // namespace cutlass
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm70, ThreadblockShape, WarpShape, GemmShape< 8, 8, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Mma
typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassTensorOp, arch::Sm70, ThreadblockShape, WarpShape, GemmShape< 8, 8, 4 >, 2, Operator >::ThreadblockMma Mma
Define the threadblock-scoped matrix multiply-accumulate.
Definition: default_gemm.h:335
cutlass::gemm::kernel::DefaultGemm
Definition: default_gemm.h:116
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Mma
typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, 2, Operator >::ThreadblockMma Mma
Define the threadblock-scoped matrix multiply-accumulate.
Definition: default_gemm.h:186
cutlass
Definition: aligned_buffer.h:35
cutlass::epilogue::threadblock::DefaultEpilogueSimt
Defines sensible defaults for epilogues for SimtOps.
Definition: default_epilogue_simt.h:70
cutlass::arch::Sm50
Definition: arch.h:37
default_epilogue_wmma_tensor_op.h
Epilogue for threadblock scoped GEMMs using Tensor Ops.
cutlass::arch::Sm70
Definition: arch.h:46
gemm.h
Defines common types used for all GEMM-like operators.
cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::Mma
typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, true >::ThreadblockMma Mma
Define the threadblock-scoped matrix multiply-accumulate.
Definition: default_gemm.h:252
default_mma.h
Template for a pipelined GEMM kernel. Does not compute batching or support split-K.
cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::Mma
typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC, arch::OpClassSimt, arch::Sm50, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, false >::ThreadblockMma Mma
Define the threadblock-scoped matrix multiply-accumulate.
Definition: default_gemm.h:502
default_mma_core_sm70.h
Defines basic properties needed by CTA-level GEMMs assuming expectations about data layout of the glo...
cutlass::gemm::threadblock::DefaultMma
Definition: default_mma.h:87
cutlass::arch::Sm75
Definition: arch.h:52
linear_combination.h
Functor performing linear combination operations used by epilogues.
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm70, ThreadblockShape, WarpShape, GemmShape< 8, 8, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Epilogue
typename cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp< ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp, EpilogueOutputOp::kCount >::Epilogue Epilogue
Define the epilogue.
Definition: default_gemm.h:346
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::gemm::GemmShape
Shape of a matrix multiply-add operation.
Definition: include/cutlass/gemm/gemm.h:57
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
default_epilogue_tensor_op.h
Epilogue for threadblock scoped GEMMs using Tensor Ops.
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 1 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Epilogue
typename cutlass::epilogue::threadblock::DefaultEpilogueSimt< ThreadblockShape, typename Mma::Operator, EpilogueOutputOp, kEpilogueElementsPerAccess >::Epilogue Epilogue
Define the epilogue.
Definition: default_gemm.h:434
default_mma_core_sm75.h
Defines basic properties needed by CTA-level GEMMs assuming expectations about data layout of the glo...
cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp
Defines sensible defaults for epilogues for TensorOps.
Definition: default_epilogue_volta_tensor_op.h:71
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Epilogue
typename cutlass::epilogue::threadblock::DefaultEpilogueTensorOp< ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp, EpilogueOutputOp::kCount >::Epilogue Epilogue
Define the epilogue.
Definition: default_gemm.h:197
cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::Epilogue
typename cutlass::epilogue::threadblock::DefaultInterleavedEpilogueTensorOp< ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp, 64/sizeof_bits< ElementC >::value, InterleavedK, IsBetaZero >::Epilogue Epilogue
Define the epilogue.
Definition: default_gemm.h:261
cutlass::layout::RowMajor
Mapping function for row-major matrices.
Definition: layout/matrix.h:50
epilogue.h
Epilogue for threadblock scoped GEMMs using Tensor Ops.
cutlass::gemm::kernel::Gemm
Definition: include/cutlass/gemm/kernel/gemm.h:52
matrix.h
Defines layout functions used by TensorRef and derived classes.
cutlass::epilogue::threadblock::DefaultEpilogueWmmaTensorOp
Defines sensible defaults for epilogues for WMMA TensorOps.
Definition: default_epilogue_wmma_tensor_op.h:71
cutlass::epilogue::threadblock::DefaultInterleavedEpilogueTensorOp
Definition: default_epilogue_tensor_op.h:147
threadblock_swizzle.h
Implements several possible threadblock-swizzling functions mapping blockIdx to GEMM problems...
cutlass::layout::ColumnMajorInterleaved
Definition: layout/matrix.h:343
cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 1 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >::Mma
typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, arch::Sm50, ThreadblockShape, WarpShape, GemmShape< 1, 1, 1 >, 2, Operator >::ThreadblockMma Mma
Define the threadblock-scoped matrix multiply-accumulate.
Definition: default_gemm.h:423
default_mma_core_simt.h
Defines basic properties needed by CTA-level GEMMs assuming expectations about data layout of the glo...
predicated_tile_iterator.h
Templates implementing loading of tiles from pitch-linear rank=2 tensors.
cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::ElementA
int8_t ElementA
Definition: default_gemm.h:481
gemm_pipelined.h
Template for a pipelined GEMM kernel. Does not compute batching or support split-K.
wmma.h
Templates exposing architecture support for warp matrix multiply-add (WMMA) operations.
cutlass::epilogue::threadblock::DefaultEpilogueTensorOp
Defines sensible defaults for epilogues for TensorOps.
Definition: default_epilogue_tensor_op.h:72
cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::Epilogue
typename cutlass::epilogue::threadblock::DefaultEpilogueSimt< ThreadblockShape, typename Mma::Operator, EpilogueOutputOp, kEpilogueElementsPerAccess >::Epilogue Epilogue
Define the epilogue.
Definition: default_gemm.h:513
cutlass.h
Basic include for CUTLASS.
gemm.h
Template for a pipelined GEMM kernel. Does not compute batching or support split-K.
default_epilogue_simt.h
Epilogue for threadblock scoped GEMMs using SIMT.
default_epilogue_volta_tensor_op.h
Epilogue for threadblock scoped GEMMs using Tensor Ops on Volta.
cutlass::layout::RowMajorInterleaved
Definition: layout/matrix.h:237