ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch > Struct Template Reference

ThreadwiseTensorSliceTransfer_v7r2&lt; SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v7r2.hpp>

Public Types

using Index = MultiIndex<nDim>
using SrcCoords = decltype(MakeCoordinates(SrcDescs{}, StaticallyIndexedArray<Index, nSrc>{}))
using DstCoords = decltype(MakeCoordinates(DstDescs{}, StaticallyIndexedArray<Index, nDst>{}))
using SrcSpaceFillingCurve
using DstSpaceFillingCurve

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v7r2 (const SrcDescs &src_descs, const StaticallyIndexedArray< Index, nSrc > &src_slice_origins, const DstDescs &dst_descs, const StaticallyIndexedArray< Index, nDst > &dst_slice_origins, const ElementwiseOperation &element_op)
template<typename Indices, enable_if_t< SrcDescs::Size()==Indices::Size(), bool > = false>
__device__ void SetSrcSliceOrigins (const SrcDescs &src_descs, const Indices &src_slice_origin_idxs)
template<typename Indices, enable_if_t< DstDescs::Size()==Indices::Size(), bool > = false>
__device__ void SetDstSliceOrigins (const DstDescs &dst_descs, const Indices &dst_slice_origin_idxs)
template<typename SrcBuffers, index_t ThreadScratchId = 0, enable_if_t< SrcDescs::Size()==SrcBuffers::Size(), bool > = false>
__device__ void RunRead (const SrcDescs &src_descs, const SrcBuffers &src_bufs, Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
template<index_t ThreadScratchId = 0>
__device__ void OOBCheck (Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
template<index_t ThreadScratchId = 0>
__device__ void TransposeFromElmToDst (Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
template<typename DstBuffers, index_t ThreadScratchId = 0, enable_if_t< DstDescs::Size()==1 &&DstBuffers::Size()==1, bool > = false>
__device__ void RunWrite (const DstDescs &dst_descs, DstBuffers dst_bufs, Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
template<typename SrcBuffers, typename DstBuffers, enable_if_t< SrcDescs::Size()==SrcBuffers::Size() &&DstDescs::Size()==DstBuffers::Size(), bool > = false>
__device__ void Run (const SrcDescs &src_descs, const SrcBuffers &src_bufs, const DstDescs &dst_descs, DstBuffers dst_bufs)
template<index_t ISrc>
__device__ void MoveSrcSliceWindow (const SrcDescs &src_descs, Number< ISrc > iSrc, const Index &src_slice_origin_step_idx)
template<index_t IDst>
__device__ void MoveDstSliceWindow (const DstDescs &dst_descs, Number< IDst > iDst, const Index &dst_slice_origin_step_idx)

Static Public Member Functions

template<typename Descs, typename Indices, enable_if_t< Descs::Size()==Indices::Size(), bool > = false>
static constexpr auto MakeCoordinates (const Descs &descs, const Indices &indices)
template<typename DataTypes, index_t ScalarPerVector>
static __device__ auto generate_vectors ()
static __device__ constexpr auto GetSrcCoordinateResetStep ()
static __device__ constexpr auto GetDstCoordinateResetStep ()
static __device__ constexpr auto GetSrcThreadScratchDescriptor ()
static __device__ constexpr auto GetDstThreadScratchDescriptor ()

Static Public Attributes

static constexpr auto I0 = Number<0>{}
static constexpr index_t nDim = SliceLengths::Size()
static constexpr index_t nSrc = SrcDescs::Size()
static constexpr index_t nDst = DstDescs::Size()
static constexpr auto src_scalar_per_access
static constexpr auto dst_scalar_per_access

Member Typedef Documentation

◆ DstCoords

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::DstCoords = decltype(MakeCoordinates(DstDescs{}, StaticallyIndexedArray<Index, nDst>{}))

◆ DstSpaceFillingCurve

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::DstSpaceFillingCurve
Initial value:
SpaceFillingCurve<SliceLengths,
DstDimAccessOrder,
remove_cv_t<decltype(dst_scalar_per_access)>,
false>
ck::remove_cv_t
typename remove_cv< T >::type remove_cv_t
Definition type.hpp:295
ck::SpaceFillingCurve
Definition tensor_space_filling_curve.hpp:20
ck::ThreadwiseTensorSliceTransfer_v7r2::dst_scalar_per_access
static constexpr auto dst_scalar_per_access
Definition threadwise_tensor_slice_transfer_v7r2.hpp:76

◆ Index

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::Index = MultiIndex<nDim>

◆ SrcCoords

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::SrcCoords = decltype(MakeCoordinates(SrcDescs{}, StaticallyIndexedArray<Index, nSrc>{}))

◆ SrcSpaceFillingCurve

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::SrcSpaceFillingCurve
Initial value:
SpaceFillingCurve<SliceLengths,
SrcDimAccessOrder,
remove_cv_t<decltype(src_scalar_per_access)>,
false>
ck::ThreadwiseTensorSliceTransfer_v7r2::src_scalar_per_access
static constexpr auto src_scalar_per_access
Definition threadwise_tensor_slice_transfer_v7r2.hpp:73

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v7r2()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::ThreadwiseTensorSliceTransfer_v7r2 ( const SrcDescs & src_descs,
const StaticallyIndexedArray< Index, nSrc > & src_slice_origins,
const DstDescs & dst_descs,
const StaticallyIndexedArray< Index, nDst > & dst_slice_origins,
const ElementwiseOperation & element_op )
inlineconstexpr

Member Function Documentation

◆ generate_vectors()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename DataTypes, index_t ScalarPerVector>
__device__ auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::generate_vectors ( )
inlinestatic

◆ GetDstCoordinateResetStep()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::GetDstCoordinateResetStep ( )
inlinestaticconstexpr

◆ GetDstThreadScratchDescriptor()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::GetDstThreadScratchDescriptor ( )
inlinestaticconstexpr

◆ GetSrcCoordinateResetStep()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::GetSrcCoordinateResetStep ( )
inlinestaticconstexpr

◆ GetSrcThreadScratchDescriptor()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::GetSrcThreadScratchDescriptor ( )
inlinestaticconstexpr

◆ MakeCoordinates()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename Descs, typename Indices, enable_if_t< Descs::Size()==Indices::Size(), bool > = false>
constexpr auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::MakeCoordinates ( const Descs & descs,
const Indices & indices )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<index_t IDst>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::MoveDstSliceWindow ( const DstDescs & dst_descs,
Number< IDst > iDst,
const Index & dst_slice_origin_step_idx )
inline

◆ MoveSrcSliceWindow()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<index_t ISrc>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::MoveSrcSliceWindow ( const SrcDescs & src_descs,
Number< ISrc > iSrc,
const Index & src_slice_origin_step_idx )
inline

◆ OOBCheck()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<index_t ThreadScratchId = 0>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::OOBCheck ( Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{})
inline

◆ Run()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename SrcBuffers, typename DstBuffers, enable_if_t< SrcDescs::Size()==SrcBuffers::Size() &&DstDescs::Size()==DstBuffers::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::Run ( const SrcDescs & src_descs,
const SrcBuffers & src_bufs,
const DstDescs & dst_descs,
DstBuffers dst_bufs )
inline

◆ RunRead()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename SrcBuffers, index_t ThreadScratchId = 0, enable_if_t< SrcDescs::Size()==SrcBuffers::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::RunRead ( const SrcDescs & src_descs,
const SrcBuffers & src_bufs,
Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{} )
inline

◆ RunWrite()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename DstBuffers, index_t ThreadScratchId = 0, enable_if_t< DstDescs::Size()==1 &&DstBuffers::Size()==1, bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::RunWrite ( const DstDescs & dst_descs,
DstBuffers dst_bufs,
Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{} )
inline

◆ SetDstSliceOrigins()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename Indices, enable_if_t< DstDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::SetDstSliceOrigins ( const DstDescs & dst_descs,
const Indices & dst_slice_origin_idxs )
inline

◆ SetSrcSliceOrigins()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<typename Indices, enable_if_t< SrcDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::SetSrcSliceOrigins ( const SrcDescs & src_descs,
const Indices & src_slice_origin_idxs )
inline

◆ TransposeFromElmToDst()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
template<index_t ThreadScratchId = 0>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::TransposeFromElmToDst ( Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{})
inline

Member Data Documentation

◆ dst_scalar_per_access

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::dst_scalar_per_access
staticconstexpr
Initial value:
= generate_sequence(
detail::lambda_scalar_per_access<DstVectorDim, DstScalarPerVector>{}, Number<nDim>{})
ck::Number
integral_constant< index_t, N > Number
Definition number.hpp:12
ck::generate_sequence
__host__ __device__ constexpr auto generate_sequence(F, Number< N >)
Definition sequence_helper.hpp:18
ck::detail::lambda_scalar_per_access
Definition threadwise_tensor_slice_transfer_util.hpp:20

◆ I0

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::I0 = Number<0>{}
staticconstexpr

◆ nDim

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::nDim = SliceLengths::Size()
staticconstexpr

◆ nDst

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::nDst = DstDescs::Size()
staticconstexpr

◆ nSrc

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::nSrc = SrcDescs::Size()
staticconstexpr

◆ src_scalar_per_access

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags, index_t NumThreadScratch = 1>
auto ck::ThreadwiseTensorSliceTransfer_v7r2< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags, NumThreadScratch >::src_scalar_per_access
staticconstexpr
Initial value:
= generate_sequence(
detail::lambda_scalar_per_access<SrcVectorDim, SrcScalarPerVector>{}, Number<nDim>{})
The documentation for this struct was generated from the following file: