CUTLASS: cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable > Class Template Reference - Cutlass

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

Public Types | Public Member Functions | Static Public Attributes | List of all members

cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable > Class Template Reference

Structure to compute the matrix product targeting CUDA cores and SIMT math instructions.

#include <mma_simt.h>

Public Types

| | using | Shape = Shape_ | | | Shape of warp-level matrix operation (concept: GemmShape) More...
| | | | using | ElementA = ElementA_ | | | Data type of multiplicand A. More...
| | | | using | LayoutA = LayoutA_ | | | Layout of multiplicand A. More...
| | | | using | ElementB = ElementB_ | | | Data type of multiplicand B. More...
| | | | using | LayoutB = LayoutB_ | | | Layout of multiplicand B. More...
| | | | using | ElementC = ElementC_ | | | Data type of accumulator matrix C. More...
| | | | using | LayoutC = LayoutC_ | | | Layout of accumulator matrix C. More...
| | | | using | Policy = Policy_ | | | Shape of the warp in units of thread (concept: MmaLanePolicySimt) More...
| | | | using | OperatorClass = arch::OpClassSimt | | | Indicates class of matrix operator. More...
| | | | using | ThreadLayoutA = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved< 4 >, LayoutA >::value, layout::ColumnMajor, typename platform::conditional< platform::is_same< layout::RowMajorInterleaved< 4 >, LayoutA >::value, layout::RowMajor, LayoutA >::type >::type | | | | using | ThreadLayoutB = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved< 4 >, LayoutB >::value, layout::ColumnMajor, typename platform::conditional< platform::is_same< layout::RowMajorInterleaved< 4 >, LayoutB >::value, layout::RowMajor, LayoutB >::type >::type | | | | using | dp4a_type = typename platform::conditional< use_dp4a, int8_t, bool >::type | | | | using | ThreadMma = thread::Mma< GemmShape< Shape::kM/Policy::WarpShape::kRow, Shape::kN/Policy::WarpShape::kColumn, Policy::LaneMmaShape::kK >, ElementA, ThreadLayoutA, ElementB, ThreadLayoutB, ElementC, LayoutC, arch::OpMultiplyAdd, dp4a_type > | | | Thread-level matrix multiply accumulate operator. More...
| | | | using | IteratorA = MmaSimtTileIterator< MatrixShape< Shape::kM, Policy::LaneMmaShape::kK >, Operand::kA, ElementA, LayoutA, Policy, PartitionsK, Shape::kK > | | | Iterates over the A operand in memory. More...
| | | | using | FragmentA = typename IteratorA::Fragment | | | Storage for A tile. More...
| | | | using | IteratorB = MmaSimtTileIterator< MatrixShape< Policy::LaneMmaShape::kK, Shape::kN >, Operand::kB, ElementB, LayoutB, Policy, PartitionsK, Shape::kK > | | | Iterates over the B operand in memory. More...
| | | | using | FragmentB = typename IteratorB::Fragment | | | Storage for B tile. More...
| | | | using | IteratorC = MmaSimtTileIterator< MatrixShape< Shape::kM, Shape::kN >, Operand::kC, ElementC, LayoutC, Policy > | | | Iterates over the C operand in memory. More...
| | | | using | FragmentC = typename ThreadMma::FragmentC | | | Storage for C tile. More...
| | |

Public Member Functions

Static Public Attributes

Member Typedef Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::dp4a_type = typename platform::conditional< use_dp4a , int8_t, bool >::type |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementA = ElementA_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementB = ElementB_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementC = ElementC_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentA = typename IteratorA::Fragment |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentB = typename IteratorB::Fragment |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentC = typename ThreadMma::FragmentC |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorA = MmaSimtTileIterator< MatrixShape<Shape::kM, Policy::LaneMmaShape::kK>, Operand::kA, ElementA, LayoutA, Policy, PartitionsK, Shape::kK > |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorB = MmaSimtTileIterator< MatrixShape<Policy::LaneMmaShape::kK, Shape::kN>, Operand::kB, ElementB, LayoutB, Policy, PartitionsK, Shape::kK > |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorC = MmaSimtTileIterator< MatrixShape<Shape::kM, Shape::kN>, Operand::kC, ElementC, LayoutC, Policy > |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutA = LayoutA_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutB = LayoutB_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutC = LayoutC_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::OperatorClass = arch::OpClassSimt |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::Policy = Policy_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::Shape = Shape_ |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadLayoutA = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutA >::value, layout::ColumnMajor, typename platform::conditional < platform::is_same< layout::RowMajorInterleaved<4>, LayoutA >::value, layout::RowMajor, LayoutA>::type >::type |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadLayoutB = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutB >::value, layout::ColumnMajor, typename platform::conditional < platform::is_same< layout::RowMajorInterleaved<4>, LayoutB >::value, layout::RowMajor, LayoutB>::type >::type |

| using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadMma = thread::Mma< GemmShape< Shape::kM / Policy::WarpShape::kRow, Shape::kN / Policy::WarpShape::kColumn, Policy::LaneMmaShape::kK>, ElementA, ThreadLayoutA, ElementB, ThreadLayoutB, ElementC, LayoutC, arch::OpMultiplyAdd, dp4a_type > |

Constructor & Destructor Documentation

| CUTLASS_DEVICE cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::MmaSimt | ( | | ) | |

| inline |

Member Function Documentation

| CUTLASS_DEVICE void cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::operator() | ( | FragmentC & | d, | | | | FragmentA const & | a, | | | | FragmentB const & | b, | | | | FragmentC const & | c, | | | | int | group_idx = 0 | | | ) | | const |

| inline |

Member Data Documentation

| constexpr bool cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::use_dp4a |

| static |

Initial value:

= (platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutA>::value ||

platform::is_same< layout::RowMajorInterleaved<4>, LayoutA >::value) &&

platform::is_same< ElementA, int8_t >::value &&

platform::is_same< ElementB, int8_t >::value

The documentation for this class was generated from the following file:

mma_simt.h

Generated by 1.8.11