Cutlass
CUDA Templates for Linear Algebra Subroutines and Solvers
Classes | Public Types | Static Public Attributes | List of all members
cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ > Struct Template Reference

#include <gemm_shared_tile.h>

Classes

struct  ThreadOffset
 Computes the thread offset in (H, W) based on thread ID. More...
 

Public Types

typedef platform::remove_const< Scalar_ >::type Scalar
 The scalar. More...
 
typedef Scalar_ * Pointer
 The pointer. More...
 
typedef Shape< kStages_, OutputTile_::kD/InstructionShape_::kD, GetExtent< kOperand, OutputTile_ >::kExtent *InstructionShape_::kD > TileWithoutSkew_
 The tile without skew. More...
 
typedef Shape< kStages_, TileWithoutSkew_::kH, TileWithoutSkew_::kW+kSkew_ > TileWithSkew
 The tile with skew. More...
 
typedef ReshapeTile< TileWithoutSkew_, kScalarsPerLds_ >::Tile TileWithoutSkew
 The tile without skew after reshaping. More...
 
typedef ReshapeTile< TileWithSkew, kScalarsPerLds_ >::Tile Tile
 The tile. More...
 
typedef Warps_ Warps
 The number of warps. More...
 
typedef ThreadsPerWarp_ ThreadsPerWarp
 The threads in a warp. More...
 
typedef Shape< 1, 1, TileWithoutSkew::kW/kWarps/kThreadsPerWarpIterations
 The number of iterations needed to load/store the tile. More...
 
typedef Shape< TileWithSkew::kW *Warps::kD, 0, kWarps *kThreadsPerWarp *kAccessSize, 0 > ImmediateOffsetStrides
 The strides in each dimension between different loads/stores. More...
 
typedef Shape< TileWithSkew::kW *Warps::kD, 0, kWarps *kThreadsPerWarp *kAccessSize, 0 > Delta
 

Static Public Attributes

static GemmOperand::Kind const kOperand = GemmOperand::kA
 
static int const kAccessSize = kScalarsPerLds_
 The number of scalars per LDG/STG. More...
 
static int const kSkew = kSkew_
 The skew. More...
 
static MemorySpace::Kind const kMemorySpace = MemorySpace::kShared
 The memory space. More...
 
static int const kWarps = GetExtent<kOperand, Warps>::kExtent
 The number of warps. More...
 
static int const kThreadsPerWarp = GetExtent<kOperand, ThreadsPerWarp>::kExtent
 The number of threads in one dimension of the warp. More...
 

Member Typedef Documentation

◆ Delta

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Shape<TileWithSkew::kW * Warps::kD, 0, kWarps * kThreadsPerWarp * kAccessSize, 0> cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Delta

◆ ImmediateOffsetStrides

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Shape<TileWithSkew::kW * Warps::kD, 0, kWarps * kThreadsPerWarp * kAccessSize, 0> cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ImmediateOffsetStrides

◆ Iterations

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Shape<1, 1, TileWithoutSkew::kW / kWarps / kThreadsPerWarp > cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Iterations

◆ Pointer

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Scalar_* cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Pointer

◆ Scalar

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef platform::remove_const<Scalar_>::type cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Scalar

◆ ThreadsPerWarp

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef ThreadsPerWarp_ cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ThreadsPerWarp

◆ Tile

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef ReshapeTile<TileWithSkew, kScalarsPerLds_>::Tile cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Tile

◆ TileWithoutSkew

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef ReshapeTile<TileWithoutSkew_, kScalarsPerLds_>::Tile cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::TileWithoutSkew

◆ TileWithoutSkew_

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Shape<kStages_, OutputTile_::kD / InstructionShape_::kD, GetExtent<kOperand, OutputTile_>::kExtent * InstructionShape_::kD> cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::TileWithoutSkew_

◆ TileWithSkew

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Shape<kStages_, TileWithoutSkew_::kH, TileWithoutSkew_::kW + kSkew_> cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::TileWithSkew

◆ Warps

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
typedef Warps_ cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::Warps

Member Data Documentation

◆ kAccessSize

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
int const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kAccessSize = kScalarsPerLds_
static

◆ kMemorySpace

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
MemorySpace::Kind const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kMemorySpace = MemorySpace::kShared
static

◆ kOperand

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
GemmOperand::Kind const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kOperand = GemmOperand::kA
static

◆ kSkew

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
int const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kSkew = kSkew_
static

◆ kThreadsPerWarp

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
int const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kThreadsPerWarp = GetExtent<kOperand, ThreadsPerWarp>::kExtent
static

◆ kWarps

template<typename Scalar_ , typename OutputTile_ , typename Warps_ , typename ThreadsPerWarp_ , typename InstructionShape_ , int kStages_, int kScalarsPerLds_, int kSkew_ = 0>
int const cutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::kWarps = GetExtent<kOperand, Warps>::kExtent
static

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