ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference

ThreadwiseTensorSliceTransfer_v6r2&lt; Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v6r2.hpp>

Public Types

using Index = MultiIndex<nDim>
using Src0Coord = decltype(make_tensor_coordinate(Src0Desc{}, Index{}))
using Src1Coord = decltype(make_tensor_coordinate(Src1Desc{}, Index{}))
using DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{}))

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v6r2 (const Src0Desc &src0_desc, const Index &src0_slice_origin, const Src1Desc &src1_desc, const Index &src1_slice_origin, const DstDesc &dst_desc, const Index &dst_slice_origin, const ElementwiseOperation &element_op)
__device__ void SetSrc0SliceOrigin (const Src0Desc &src0_desc, const Index &src0_slice_origin_idx)
__device__ void SetSrc1SliceOrigin (const Src1Desc &src1_desc, const Index &src1_slice_origin_idx)
__device__ void SetDstSliceOrigin (const DstDesc &dst_desc, const Index &dst_slice_origin_idx)
template<typename Src0Buffer, typename Src1Buffer, typename DstBuffer>
__device__ void Run (const Src0Desc &src0_desc, const Src0Buffer &src0_buf, const Src1Desc &src1_desc, const Src1Buffer &src1_buf, const DstDesc &dst_desc, DstBuffer &dst_buf)
__device__ void MoveSrc0SliceWindow (const Src0Desc &src0_desc, const Index &src0_slice_origin_step_idx)
__device__ void MoveSrc1SliceWindow (const Src1Desc &src1_desc, const Index &src1_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 Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{}))

◆ Index

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Index = MultiIndex<nDim>

◆ Src0Coord

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Src0Coord = decltype(make_tensor_coordinate(Src0Desc{}, Index{}))

◆ Src1Coord

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
using ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Src1Coord = decltype(make_tensor_coordinate(Src1Desc{}, Index{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v6r2()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::ThreadwiseTensorSliceTransfer_v6r2 ( const Src0Desc & src0_desc,
const Index & src0_slice_origin,
const Src1Desc & src1_desc,
const Index & src1_slice_origin,
const DstDesc & dst_desc,
const Index & dst_slice_origin,
const ElementwiseOperation & element_op )
inlineconstexpr

Member Function Documentation

◆ GetCoordinateResetStep()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::GetCoordinateResetStep ( )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveDstSliceWindow ( const DstDesc & dst_desc,
const Index & dst_slice_origin_step_idx )
inline

◆ MoveSrc0SliceWindow()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrc0SliceWindow ( const Src0Desc & src0_desc,
const Index & src0_slice_origin_step_idx )
inline

◆ MoveSrc1SliceWindow()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrc1SliceWindow ( const Src1Desc & src1_desc,
const Index & src1_slice_origin_step_idx )
inline

◆ Run()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
template<typename Src0Buffer, typename Src1Buffer, typename DstBuffer>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Run ( const Src0Desc & src0_desc,
const Src0Buffer & src0_buf,
const Src1Desc & src1_desc,
const Src1Buffer & src1_buf,
const DstDesc & dst_desc,
DstBuffer & dst_buf )
inline

◆ SetDstSliceOrigin()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetDstSliceOrigin ( const DstDesc & dst_desc,
const Index & dst_slice_origin_idx )
inline

◆ SetSrc0SliceOrigin()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrc0SliceOrigin ( const Src0Desc & src0_desc,
const Index & src0_slice_origin_idx )
inline

◆ SetSrc1SliceOrigin()

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
__device__ void ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrc1SliceOrigin ( const Src1Desc & src1_desc,
const Index & src1_slice_origin_idx )
inline

Member Data Documentation

◆ I0

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
auto ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::I0 = Number<0>{}
staticconstexpr

◆ nDim

template<typename Src0Data, typename Src1Data, typename DstData, typename Src0Desc, typename Src1Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
index_t ck::ThreadwiseTensorSliceTransfer_v6r2< Src0Data, Src1Data, DstData, Src0Desc, Src1Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::nDim = SliceLengths::Size()
staticconstexpr
The documentation for this struct was generated from the following file: