CUTLASS: cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess > Struct Template Reference

CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers

#include <pitch_linear_thread_map.h>

Classes
struct	Detail
	Internal details made public to facilitate introspection Iterations along each dimension (concept: PitchLinearShape) More...

Public Types
using	TensorCoord = layout::PitchLinearCoord
	Tensor coordinate. More...

using	Shape = Shape_
	Tile shape. More...

using	ThreadAccessShape = layout::PitchLinearShape< kElementsPerAccess, 1 >
	Shape of access by each thread. More...

using	Iterations = layout::PitchLinearShape< Detail::WarpAccessIterations::kContiguous/Detail::kWarpsContiguous, Detail::WarpAccessIterations::kStrided/Detail::kWarpsStrided >

using	Delta = layout::PitchLinearShape< Detail::WarpThreadArrangement::kContiguous kElementsPerAccess, Detail::WarpThreadArrangement::kStrided Detail::WarpArrangement::kStrided >
	Delta betweeen accesses (units of elements, concept: PitchLinearShape) More...

Static Public Member Functions
static CUTLASS_HOST_DEVICE TensorCoord	initial_offset (int thread_id)
	Maps thread ID to a coordinate offset within the tensor's logical coordinate space. More...

Static Public Attributes
static int const	kThreads = Threads
	Number of threads total. More...

static int const	kElementsPerAccess = ElementsPerAccess
	Extract vector length from Layout. More...

Detailed Description

template<typename Shape_, int Threads, typename WarpThreadArrangement_, int ElementsPerAccess = 1>
struct cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >

Policy defining a warp-striped arrangement. This partitions a tile into vectorized memory accesses performed by each warp then distributes warps across them. Warps are striped in the strided dimension and raked across the contiguous dimension.

Member Typedef Documentation

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

using cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::Delta = layout::PitchLinearShape< Detail::WarpThreadArrangement::kContiguous * kElementsPerAccess, Detail::WarpThreadArrangement::kStrided * Detail::WarpArrangement::kStrided >

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

using cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::Iterations = layout::PitchLinearShape< Detail::WarpAccessIterations::kContiguous / Detail::kWarpsContiguous, Detail::WarpAccessIterations::kStrided / Detail::kWarpsStrided >

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

using cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::Shape = Shape_

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

using cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::TensorCoord = layout::PitchLinearCoord

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

using cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::ThreadAccessShape = layout::PitchLinearShape<kElementsPerAccess, 1>

Member Function Documentation

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

static CUTLASS_HOST_DEVICE TensorCoord cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::initial_offset ( int thread_id )

inlinestatic

Member Data Documentation

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

int const cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::kElementsPerAccess = ElementsPerAccess

static

template<typename Shape_ , int Threads, typename WarpThreadArrangement_ , int ElementsPerAccess = 1>

int const cutlass::transform::PitchLinearWarpStripedThreadMap< Shape_, Threads, WarpThreadArrangement_, ElementsPerAccess >::kThreads = Threads

static

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

pitch_linear_thread_map.h