ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference

ThreadwiseTensorSliceTransfer_v6r1&lt; SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v6r1.hpp>

Public Types

using Index = MultiIndex<nDim>
using SrcCoord = decltype(make_tensor_coordinate(SrcDesc{}, Index{}))
using DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{}))

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v6r1 (const SrcDesc &src_desc, const Index &src_slice_origin, const DstDesc &dst_desc, const Index &dst_slice_origin, const ElementwiseOperation &element_op)
__device__ void SetSrcSliceOrigin (const SrcDesc &src_desc, const Index &src_slice_origin_idx)
__device__ void SetDstSliceOrigin (const DstDesc &dst_desc, const Index &dst_slice_origin_idx)
template<typename SrcBuffer, typename DstBuffer>
__device__ void Run (const SrcDesc &src_desc, const SrcBuffer &src_buf, const DstDesc &dst_desc, DstBuffer &dst_buf)
__device__ void MoveSrcSliceWindow (const SrcDesc &src_desc, const Index &src_slice_origin_step_idx)
__device__ void MoveDstSliceWindow (const DstDesc &dst_desc, const Index &dst_slice_origin_step_idx)

Static Public Member Functions

static __device__ constexpr auto GetCoordinateResetStep ()

Static Public Attributes

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

Member Typedef Documentation

◆ DstCoord

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{}))

◆ Index

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Index = MultiIndex<nDim>

◆ SrcCoord

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SrcCoord = decltype(make_tensor_coordinate(SrcDesc{}, Index{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v6r1()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::ThreadwiseTensorSliceTransfer_v6r1 ( const SrcDesc & src_desc,
const Index & src_slice_origin,
const DstDesc & dst_desc,
const Index & dst_slice_origin,
const ElementwiseOperation & element_op )
inlineconstexpr

Member Function Documentation

◆ GetCoordinateResetStep()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::GetCoordinateResetStep ( )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveDstSliceWindow ( const DstDesc & dst_desc,
const Index & dst_slice_origin_step_idx )
inline

◆ MoveSrcSliceWindow()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrcSliceWindow ( const SrcDesc & src_desc,
const Index & src_slice_origin_step_idx )
inline

◆ Run()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
template<typename SrcBuffer, typename DstBuffer>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Run ( const SrcDesc & src_desc,
const SrcBuffer & src_buf,
const DstDesc & dst_desc,
DstBuffer & dst_buf )
inline

◆ SetDstSliceOrigin()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetDstSliceOrigin ( const DstDesc & dst_desc,
const Index & dst_slice_origin_idx )
inline

◆ SetSrcSliceOrigin()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrcSliceOrigin ( const SrcDesc & src_desc,
const Index & src_slice_origin_idx )
inline

Member Data Documentation

◆ I0

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
auto ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::I0 = Number<0>{}
staticconstexpr

◆ nDim

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool SrcResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
index_t ck::ThreadwiseTensorSliceTransfer_v6r1< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, SrcResetCoordinateAfterRun, DstResetCoordinateAfterRun >::nDim = SliceLengths::Size()
staticconstexpr
The documentation for this struct was generated from the following file: