Back to Cutlass

CUTLASS: default_epilogue_volta_tensor_op.h Source File

docs/default__epilogue__volta__tensor__op_8h_source.html

4.4.221.4 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

default_epilogue_volta_tensor_op.h

[Go to the documentation of this file.](default epilogue volta tensor op_8h.html)

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 **************************************************************************************************/

33 #pragma once

34

35 #include "cutlass/cutlass.h"

36 #include "cutlass/numeric_types.h"

37 #include "cutlass/array.h"

38

39 #include "cutlass/gemm/gemm.h"

40

41 #include "cutlass/epilogue/thread/linear_combination.h"

42 #include "cutlass/epilogue/thread/conversion_op.h"

43 #include "cutlass/epilogue/thread/reduction_op.h"

44

45 #include "[cutlass/transform/threadblock/regular_tile_iterator_pitch_linear.h](regular tile iterator pitch linear_8h.html)"

46 #include "[cutlass/epilogue/threadblock/predicated_tile_iterator.h](epilogue_2threadblock_2predicated tile iterator_8h.html)"

47 #include "[cutlass/epilogue/threadblock/shared_load_iterator.h](shared load iterator_8h.html)"

48

49 #include "[cutlass/epilogue/warp/fragment_iterator_volta_tensor_op.h](fragment iterator volta tensor op_8h.html)"

50 #include "[cutlass/epilogue/warp/tile_iterator_volta_tensor_op.h](tile iterator volta tensor op_8h.html)"

51 #include "[cutlass/epilogue/threadblock/default_thread_map_volta_tensor_op.h](default thread map volta tensor__op_8h.html)"

52

53 #include "cutlass/epilogue/threadblock/epilogue.h"

54

56

57 namespace cutlass {

58 namespace epilogue {

59 namespace threadblock {

60

62

64 template <

65typename Shape_,

66typename WarpMmaTensorOp_,

67int PartitionsK,

68typename OutputOp_,

69int ElementsPerAccess

70 >

71 struct DefaultEpilogueVoltaTensorOp {

72

73using Shape = Shape_;

74using WarpMmaTensorOp = WarpMmaTensorOp_;

75static int const kPartitionsK = PartitionsK;

76using OutputOp = OutputOp_;

77static int const kElementsPerAccess = ElementsPerAccess;

78

79using ElementOutput = typename OutputOp::ElementOutput;

80using LayoutC = typename WarpMmaTensorOp::LayoutC;

81using ElementAccumulator = typename WarpMmaTensorOp::ElementC;

82

83//

84// Thread map

85//

86

87using OutputTileThreadMap = typename cutlass::epilogue::threadblock::DefaultThreadMapVoltaTensorOp<

88Shape,

89typename WarpMmaTensorOp::Shape,

90kPartitionsK,

91ElementOutput,

92kElementsPerAccess,

93ElementAccumulator

94 >::Type;

95

96using OutputTileIterator = cutlass::epilogue::threadblock::PredicatedTileIterator<

97OutputTileThreadMap,

98 ElementOutput

99 >;

100

101using AccumulatorFragmentIterator = cutlass::epilogue::warp::FragmentIteratorVoltaTensorOp<

102typename WarpMmaTensorOp::Shape,

103gemm::GemmShape<32, 32, 4>,

104ElementAccumulator,

105LayoutC

106 >;

107

108using WarpTileIterator = cutlass::epilogue::warp::TileIteratorVoltaTensorOp<

109typename WarpMmaTensorOp::Shape,

110 gemm::GemmShape<32, 32, 4>,

111ElementAccumulator,

112LayoutC

113 >;

114

115static int const kSharedMemAlignment = sizeof_bits<ElementAccumulator>::value * WarpTileIterator::kElementsPerAccess / 8;

116

117static_assert(kSharedMemAlignment == 8, "Shared memory alignment must be 8B");

118

119using SharedLoadIterator = cutlass::epilogue::threadblock::SharedLoadIterator<

120typename OutputTileThreadMap::CompactedThreadMap,

121ElementAccumulator,

122 kSharedMemAlignment

123 >;

124

126using Padding = typename WarpTileIterator::Padding;

127

128//

129// Define the epilogue

130//

131using Epilogue = cutlass::epilogue::threadblock::Epilogue<

132Shape,

133WarpMmaTensorOp,

134kPartitionsK,

135OutputTileIterator,

136AccumulatorFragmentIterator,

137WarpTileIterator,

138SharedLoadIterator,

139OutputOp,

140Padding

141 >;

142 };

143

145

146 } // namespace threadblock

147 } // namespace epilogue

148 } // namespace cutlass

149

[regular_tile_iterator_pitch_linear.h](regular tile iterator pitch linear_8h.html)

Templates implementing loading of tiles from pitch-linear rank=2 tensors.

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::WarpMmaTensorOp

WarpMmaTensorOp_ WarpMmaTensorOp

Definition: default_epilogue_volta_tensor_op.h:74

cutlass

Definition: aligned_buffer.h:35

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::kPartitionsK

static int const kPartitionsK

Definition: default_epilogue_volta_tensor_op.h:75

[predicated_tile_iterator.h](epilogue_2threadblock_2predicated tile iterator_8h.html)

Epilogue for threadblock scoped GEMMs using Tensor Ops.

[fragment_iterator_volta_tensor_op.h](fragment iterator volta tensor op_8h.html)

This defines a "fragment" iterator for visiting the fragments of an accumulator tile that participate...

gemm.h

Defines common types used for all GEMM-like operators.

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::kElementsPerAccess

static int const kElementsPerAccess

Definition: default_epilogue_volta_tensor_op.h:77

conversion_op.h

Functor performing conversion operations used by epilogues.

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::LayoutC

typename WarpMmaTensorOp::LayoutC LayoutC

Definition: default_epilogue_volta_tensor_op.h:80

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::Shape

Shape_ Shape

Definition: default_epilogue_volta_tensor_op.h:73

array.h

Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::SharedLoadIterator

cutlass::epilogue::threadblock::SharedLoadIterator< typename OutputTileThreadMap::CompactedThreadMap, ElementAccumulator, kSharedMemAlignment > SharedLoadIterator

Definition: default_epilogue_volta_tensor_op.h:123

[tile_iterator_volta_tensor_op.h](tile iterator volta tensor op_8h.html)

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::epilogue::warp::TileIteratorVoltaTensorOp

Template for reading and writing tiles of accumulators to shared memory.

Definition: tile_iterator_volta_tensor_op.h:52

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::AccumulatorFragmentIterator

cutlass::epilogue::warp::FragmentIteratorVoltaTensorOp< typename WarpMmaTensorOp::Shape, gemm::GemmShape< 32, 32, 4 >, ElementAccumulator, LayoutC > AccumulatorFragmentIterator

Definition: default_epilogue_volta_tensor_op.h:106

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::kSharedMemAlignment

static int const kSharedMemAlignment

Definition: default_epilogue_volta_tensor_op.h:115

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::WarpTileIterator

cutlass::epilogue::warp::TileIteratorVoltaTensorOp< typename WarpMmaTensorOp::Shape, gemm::GemmShape< 32, 32, 4 >, ElementAccumulator, LayoutC > WarpTileIterator

Definition: default_epilogue_volta_tensor_op.h:113

[default_thread_map_volta_tensor_op.h](default thread map volta tensor__op_8h.html)

numeric_types.h

Top-level include for all CUTLASS numeric types.

cutlass::epilogue::warp::FragmentIteratorVoltaTensorOp

Definition: fragment_iterator_volta_tensor_op.h:61

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

[shared_load_iterator.h](shared load iterator_8h.html)

Epilogue for threadblock scoped GEMMs using Tensor Ops.

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp

Defines sensible defaults for epilogues for TensorOps.

Definition: default_epilogue_volta_tensor_op.h:71

cutlass::epilogue::threadblock::Epilogue

Epilogue operator without splitk.

Definition: epilogue.h:74

epilogue.h

Epilogue for threadblock scoped GEMMs using Tensor Ops.

cutlass::epilogue::threadblock::PredicatedTileIterator

Definition: epilogue/threadblock/predicated_tile_iterator.h:65

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::OutputTileIterator

cutlass::epilogue::threadblock::PredicatedTileIterator< OutputTileThreadMap, ElementOutput > OutputTileIterator

Definition: default_epilogue_volta_tensor_op.h:99

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::ElementAccumulator

typename WarpMmaTensorOp::ElementC ElementAccumulator

Definition: default_epilogue_volta_tensor_op.h:81

cutlass::epilogue::threadblock::DefaultThreadMapVoltaTensorOp

Defines the optimal thread map for TensorOp accumulator layouts.

Definition: default_thread_map_volta_tensor_op.h:52

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::Padding

typename WarpTileIterator::Padding Padding

Hard-coded padding elements added.

Definition: default_epilogue_volta_tensor_op.h:126

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::OutputTileThreadMap

typename cutlass::epilogue::threadblock::DefaultThreadMapVoltaTensorOp< Shape, typename WarpMmaTensorOp::Shape, kPartitionsK, ElementOutput, kElementsPerAccess, ElementAccumulator >::Type OutputTileThreadMap

Definition: default_epilogue_volta_tensor_op.h:94

cutlass::epilogue::threadblock::SharedLoadIterator

Definition: shared_load_iterator.h:61

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::ElementOutput

typename OutputOp::ElementOutput ElementOutput

Definition: default_epilogue_volta_tensor_op.h:79

reduction_op.h

Functor performing reduction operations used by epilogues.

cutlass.h

Basic include for CUTLASS.

cutlass::epilogue::threadblock::DefaultEpilogueVoltaTensorOp::OutputOp

OutputOp_ OutputOp

Definition: default_epilogue_volta_tensor_op.h:76


Generated by 1.8.11