ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type > Struct Template Reference

ThreadwiseTensorSliceTransfer_v1r3&lt; SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type > Struct Template Reference

#include <threadwise_tensor_slice_transfer.hpp>

Public Types

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

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v1r3 (const DstDesc &dst_desc, const Index &dst_slice_origin_idx, const ElementwiseOperation &element_op)
__device__ void SetDstSliceOrigin (const DstDesc &dst_desc, const Index &dst_slice_origin_idx)
template<typename SrcSliceOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void Run (const SrcDesc &, const SrcSliceOriginIdx &, const SrcBuffer &src_buf, const DstDesc &dst_desc, DstBuffer &dst_buf)
__device__ void MoveDstSliceWindow (const DstDesc &dst_desc, const Index &dst_slice_origin_step_idx)

Static Public Member Functions

static __device__ constexpr auto GetDstCoordinateResetStep ()

Static Public Attributes

static constexpr index_t nDim = SliceLengths::Size()

Member Typedef Documentation

◆ DstCoord

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{}))

◆ DstCoordStep

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::DstCoordStep = decltype(make_tensor_coordinate_step(DstDesc{}, Index{}))

◆ Index

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::Index = MultiIndex<nDim>

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v1r3()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::ThreadwiseTensorSliceTransfer_v1r3 ( const DstDesc & dst_desc,
const Index & dst_slice_origin_idx,
const ElementwiseOperation & element_op )
inlineconstexpr

Member Function Documentation

◆ GetDstCoordinateResetStep()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::GetDstCoordinateResetStep ( )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::MoveDstSliceWindow ( const DstDesc & dst_desc,
const Index & dst_slice_origin_step_idx )
inline

◆ Run()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcSliceOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::Run ( const SrcDesc & ,
const SrcSliceOriginIdx & ,
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 DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::SetDstSliceOrigin ( const DstDesc & dst_desc,
const Index & dst_slice_origin_idx )
inline

Member Data Documentation

◆ nDim

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, InMemoryDataOperationEnum DstInMemOp, index_t DstScalarStrideInVector, bool DstResetCoordinateAfterRun, typename enable_if< SrcDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v1r3< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, DstInMemOp, DstScalarStrideInVector, DstResetCoordinateAfterRun, type >::nDim = SliceLengths::Size()
staticconstexpr
The documentation for this struct was generated from the following file: