CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable > Class Template Reference

Partial specialization for complex*complex+complex => complex using real-valued TensorOps.

#include <mma_complex_tensor_op.h>

Public Types

using Shape = Shape_
 Shape of warp-level matrix operation (concept: GemmShape) More...
 
using ElementA = complex< RealElementA >
 Data type of multiplicand A. More...
 
using LayoutA = LayoutA_
 Layout of multiplicand A. More...
 
using ElementB = complex< RealElementB >
 Data type of multiplicand B. More...
 
using LayoutB = LayoutB_
 Layout of multiplicand B. More...
 
using ElementC = complex< RealElementC >
 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::OpClassTensorOp
 Indicates class of matrix operator. More...
 
using IteratorA = MmaTensorOpMultiplicandTileIterator< MatrixShape< Shape::kM, Shape::kK >, Operand::kA, ElementA, LayoutA, MatrixShape< Policy::Operator::Shape::kM, Policy::Operator::Shape::kK >, Policy::OpDelta::kRow, 32, 1 >
 Iterates over the A operand in memory. More...
 
using FragmentA = typename IteratorA::Fragment
 Storage for A tile. More...
 
using IteratorB = MmaTensorOpMultiplicandTileIterator< MatrixShape< Shape::kK, Shape::kN >, Operand::kB, ElementB, LayoutB, MatrixShape< Policy::Operator::Shape::kK, Policy::Operator::Shape::kN >, Policy::OpDelta::kColumn, 32, 1 >
 Iterates over the B operand in memory. More...
 
using FragmentB = typename IteratorB::Fragment
 Storage for B tile. More...
 
using MmaIterations = MatrixShape< Shape::kM/Policy::Operator::Shape::kM, Shape::kN/Policy::Operator::Shape::kN >
 Number of mma operations performed. More...
 
using IteratorC = MmaTensorOpAccumulatorTileIterator< MatrixShape< Shape::kM, Shape::kN >, ElementC, LayoutC, typename Policy::Operator::Shape, typename Policy::OpDelta >
 Iterates over the C operand in memory. More...
 
using FragmentC = typename IteratorC::Fragment
 

Public Member Functions

CUTLASS_DEVICE MmaComplexTensorOp ()
 Ctor. More...
 
CUTLASS_DEVICE void operator() (FragmentC &D, FragmentA const &A, FragmentB const &B, FragmentC const &C) const
 Performs a warp-level matrix multiply-accumulate operation. More...
 

Static Public Attributes

static ComplexTransform const kTransformA = TransformA
 Complex transform on A operand. More...
 
static ComplexTransform const kTransformB = TransformB
 Complex transform on B operand. More...
 
static int const kThreadCount = 32
 Number of threads participating in warp-level matrix product. More...
 

Member Typedef Documentation

template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::ElementA = complex<RealElementA>
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::ElementB = complex<RealElementB>
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::ElementC = complex<RealElementC>
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::FragmentA = typename IteratorA::Fragment
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::FragmentB = typename IteratorB::Fragment
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::FragmentC = typename IteratorC::Fragment

Storage for C tile, the accumulator. Note, regardless of multiplicand type, this storage arrangement is to be considered 'planar complex' in the sense that all real-valued parts are stored consecutively followed by all imaginary parts. This matches the structure of Tensor Cores which are always real-valued matrix multiplies.

template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::IteratorA = MmaTensorOpMultiplicandTileIterator< MatrixShape<Shape::kM, Shape::kK>, Operand::kA, ElementA, LayoutA, MatrixShape<Policy::Operator::Shape::kM, Policy::Operator::Shape::kK>, Policy::OpDelta::kRow, 32, 1 >
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::IteratorB = MmaTensorOpMultiplicandTileIterator< MatrixShape<Shape::kK, Shape::kN>, Operand::kB, ElementB, LayoutB, MatrixShape<Policy::Operator::Shape::kK, Policy::Operator::Shape::kN>, Policy::OpDelta::kColumn, 32, 1 >
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::IteratorC = MmaTensorOpAccumulatorTileIterator< MatrixShape<Shape::kM, Shape::kN>, ElementC, LayoutC, typename Policy::Operator::Shape, typename Policy::OpDelta>
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::LayoutA = LayoutA_
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::LayoutB = LayoutB_
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::LayoutC = LayoutC_
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::MmaIterations = MatrixShape< Shape::kM / Policy::Operator::Shape::kM, Shape::kN / Policy::Operator::Shape::kN >
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::OperatorClass = arch::OpClassTensorOp
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::Policy = Policy_
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
using cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::Shape = Shape_

Constructor & Destructor Documentation

template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
CUTLASS_DEVICE cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::MmaComplexTensorOp ( )
inline

Member Function Documentation

template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
CUTLASS_DEVICE void cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::operator() ( FragmentC &  D,
FragmentA const &  A,
FragmentB const &  B,
FragmentC const &  C 
) const
inline

Member Data Documentation

template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
int const cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::kThreadCount = 32
static
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
ComplexTransform const cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::kTransformA = TransformA
static
template<typename Shape_ , typename RealElementA , typename LayoutA_ , typename RealElementB , typename LayoutB_ , typename RealElementC , typename LayoutC_ , typename Policy_ , ComplexTransform TransformA, ComplexTransform TransformB, typename Enable >
ComplexTransform const cutlass::gemm::warp::MmaComplexTensorOp< Shape_, complex< RealElementA >, LayoutA_, complex< RealElementB >, LayoutB_, complex< RealElementC >, LayoutC_, Policy_, TransformA, TransformB, Enable >::kTransformB = TransformB
static
The documentation for this class was generated from the following file:
Generated by   doxygen 1.8.11