ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch > Struct Template Reference

ThreadwiseTensorSliceTransfer_v3r2&lt; SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v3r2.hpp>

Public Types

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

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v3r2 (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>
__device__ void RunRead (const SrcDescs &src_descs, const SrcBuffers &src_bufs, Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
template<index_t ThreadScratchId>
__device__ void TransferDataFromSrcThreadScratchToDstThreadScratch (Number< ThreadScratchId > thread_scratch_id)
template<typename DstBuffers, index_t ThreadScratchId = 0>
__device__ void RunWrite (const DstDescs &dst_descs, DstBuffers &dst_bufs, Number< ThreadScratchId > thread_scratch_id=Number< ThreadScratchId >{})
__device__ void MoveSrcSliceWindow (const SrcDescs &src_descs, const Index &src_slice_origin_step_idx)
__device__ void MoveDstSliceWindow (const DstDescs &dst_descs, 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<index_t src_i>
static __device__ constexpr auto GetSrcCoordinateResetStep ()
template<index_t dst_i>
static __device__ constexpr auto GetDstCoordinateResetStep ()
template<index_t src_i>
static __device__ constexpr auto GetSrcThreadScratchDescriptor ()
template<index_t dst_i>
static __device__ constexpr auto GetDstThreadScratchDescriptor ()
static __device__ constexpr auto MakeSrcThreadScratchTuple ()
static __device__ constexpr auto MakeDstThreadScratchTuple ()

Static Public Attributes

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

Member Typedef Documentation

◆ DstCoords

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::DstCoords = decltype(MakeCoordinates(DstDescs{}, StaticallyIndexedArray<Index, nDst>{}))

◆ Index

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::Index = MultiIndex<nDim>

◆ SrcCoords

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
using ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::SrcCoords = decltype(MakeCoordinates(SrcDescs{}, StaticallyIndexedArray<Index, nSrc>{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v3r2()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::ThreadwiseTensorSliceTransfer_v3r2 ( 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

◆ GetDstCoordinateResetStep()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<index_t dst_i>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::GetDstCoordinateResetStep ( )
inlinestaticconstexpr

◆ GetDstThreadScratchDescriptor()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<index_t dst_i>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::GetDstThreadScratchDescriptor ( )
inlinestaticconstexpr

◆ GetSrcCoordinateResetStep()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<index_t src_i>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::GetSrcCoordinateResetStep ( )
inlinestaticconstexpr

◆ GetSrcThreadScratchDescriptor()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<index_t src_i>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::GetSrcThreadScratchDescriptor ( )
inlinestaticconstexpr

◆ MakeCoordinates()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<typename Descs, typename Indices, enable_if_t< Descs::Size()==Indices::Size(), bool > = false>
constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::MakeCoordinates ( const Descs & descs,
const Indices & indices )
inlinestaticconstexpr

◆ MakeDstThreadScratchTuple()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::MakeDstThreadScratchTuple ( )
inlinestaticconstexpr

◆ MakeSrcThreadScratchTuple()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::MakeSrcThreadScratchTuple ( )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::MoveDstSliceWindow ( const DstDescs & dst_descs,
const Index & dst_slice_origin_step_idx )
inline

◆ MoveSrcSliceWindow()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::MoveSrcSliceWindow ( const SrcDescs & src_descs,
const Index & src_slice_origin_step_idx )
inline

◆ RunRead()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<typename SrcBuffers, index_t ThreadScratchId = 0>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::RunRead ( const SrcDescs & src_descs,
const SrcBuffers & src_bufs,
Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{} )
inline

◆ RunWrite()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<typename DstBuffers, index_t ThreadScratchId = 0>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::RunWrite ( const DstDescs & dst_descs,
DstBuffers & dst_bufs,
Number< ThreadScratchId > thread_scratch_id = Number<ThreadScratchId>{} )
inline

◆ SetDstSliceOrigins()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<typename Indices, enable_if_t< DstDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::SetDstSliceOrigins ( const DstDescs & dst_descs,
const Indices & dst_slice_origin_idxs )
inline

◆ SetSrcSliceOrigins()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<typename Indices, enable_if_t< SrcDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::SetSrcSliceOrigins ( const SrcDescs & src_descs,
const Indices & src_slice_origin_idxs )
inline

◆ TransferDataFromSrcThreadScratchToDstThreadScratch()

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
template<index_t ThreadScratchId>
__device__ void ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::TransferDataFromSrcThreadScratchToDstThreadScratch ( Number< ThreadScratchId > thread_scratch_id)
inline

Member Data Documentation

◆ I0

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
auto ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::I0 = Number<0>{}
staticconstexpr

◆ nDim

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::nDim = SliceLengths::Size()
staticconstexpr

◆ nDst

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::nDst = DstDescs::Size()
staticconstexpr

◆ nSrc

template<typename SliceLengths, typename ElementwiseOperation, typename DstInMemOps, typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename SrcDimAccessOrder, typename DstDimAccessOrder, index_t SrcVectorDim, index_t DstVectorDim, typename SrcsScalarPerVector, typename DstsScalarPerVector, typename SrcsScalarStrideInVector, typename DstsScalarStrideInVector, typename SrcsResetCoordinateAfterRun, typename DstsResetCoordinateAfterRun, index_t NumThreadScratch = 1>
index_t ck::ThreadwiseTensorSliceTransfer_v3r2< SliceLengths, ElementwiseOperation, DstInMemOps, SrcDatas, DstDatas, SrcDescs, DstDescs, SrcDimAccessOrder, DstDimAccessOrder, SrcVectorDim, DstVectorDim, SrcsScalarPerVector, DstsScalarPerVector, SrcsScalarStrideInVector, DstsScalarStrideInVector, SrcsResetCoordinateAfterRun, DstsResetCoordinateAfterRun, NumThreadScratch >::nSrc = SrcDescs::Size()
staticconstexpr
The documentation for this struct was generated from the following file: