#include <threadwise_tensor_slice_transfer_v6r3.hpp>
|
| __device__ constexpr | ThreadwiseTensorSliceTransfer_v6r3 (const Src0Desc &src0_desc, const Index &src0_slice_origin, const Src1Desc &src1_desc, const Index &src1_slice_origin, const Src2Desc &src2_desc, const Index &src2_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 | SetSrc2SliceOrigin (const Src2Desc &src2_desc, const Index &src2_slice_origin_idx) |
| __device__ void | SetDstSliceOrigin (const DstDesc &dst_desc, const Index &dst_slice_origin_idx) |
| template<typename Src0Buffer, typename Src1Buffer, typename Src2Buffer, typename DstBuffer> |
| __device__ void | Run (const Src0Desc &src0_desc, const Src0Buffer &src0_buf, const Src1Desc &src1_desc, const Src1Buffer &src1_buf, const Src2Desc &src2_desc, const Src2Buffer &src2_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 | MoveSrc2SliceWindow (const Src2Desc &src2_desc, const Index &src2_slice_origin_step_idx) |
| __device__ void | MoveDstSliceWindow (const DstDesc &dst_desc, const Index &dst_slice_origin_step_idx) |
◆ DstCoord
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| using ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::DstCoord = decltype(make_tensor_coordinate(DstDesc{}, Index{})) |
◆ Index
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| using ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Index = MultiIndex<nDim> |
◆ Src0Coord
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| using ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Src0Coord = decltype(make_tensor_coordinate(Src0Desc{}, Index{})) |
◆ Src1Coord
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| using ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Src1Coord = decltype(make_tensor_coordinate(Src1Desc{}, Index{})) |
◆ Src2Coord
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| using ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Src2Coord = decltype(make_tensor_coordinate(Src2Desc{}, Index{})) |
◆ ThreadwiseTensorSliceTransfer_v6r3()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ constexpr ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::ThreadwiseTensorSliceTransfer_v6r3 |
( |
const Src0Desc & | src0_desc, |
|
|
const Index & | src0_slice_origin, |
|
|
const Src1Desc & | src1_desc, |
|
|
const Index & | src1_slice_origin, |
|
|
const Src2Desc & | src2_desc, |
|
|
const Index & | src2_slice_origin, |
|
|
const DstDesc & | dst_desc, |
|
|
const Index & | dst_slice_origin, |
|
|
const ElementwiseOperation & | element_op ) |
|
inlineconstexpr |
◆ GetCoordinateResetStep()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::GetCoordinateResetStep |
( |
| ) |
|
|
inlinestaticconstexpr |
◆ MoveDstSliceWindow()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveDstSliceWindow |
( |
const DstDesc & | dst_desc, |
|
|
const Index & | dst_slice_origin_step_idx ) |
|
inline |
◆ MoveSrc0SliceWindow()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrc0SliceWindow |
( |
const Src0Desc & | src0_desc, |
|
|
const Index & | src0_slice_origin_step_idx ) |
|
inline |
◆ MoveSrc1SliceWindow()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrc1SliceWindow |
( |
const Src1Desc & | src1_desc, |
|
|
const Index & | src1_slice_origin_step_idx ) |
|
inline |
◆ MoveSrc2SliceWindow()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::MoveSrc2SliceWindow |
( |
const Src2Desc & | src2_desc, |
|
|
const Index & | src2_slice_origin_step_idx ) |
|
inline |
◆ Run()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
template<typename Src0Buffer, typename Src1Buffer, typename Src2Buffer, typename DstBuffer>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::Run |
( |
const Src0Desc & | src0_desc, |
|
|
const Src0Buffer & | src0_buf, |
|
|
const Src1Desc & | src1_desc, |
|
|
const Src1Buffer & | src1_buf, |
|
|
const Src2Desc & | src2_desc, |
|
|
const Src2Buffer & | src2_buf, |
|
|
const DstDesc & | dst_desc, |
|
|
DstBuffer & | dst_buf ) |
|
inline |
◆ SetDstSliceOrigin()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetDstSliceOrigin |
( |
const DstDesc & | dst_desc, |
|
|
const Index & | dst_slice_origin_idx ) |
|
inline |
◆ SetSrc0SliceOrigin()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrc0SliceOrigin |
( |
const Src0Desc & | src0_desc, |
|
|
const Index & | src0_slice_origin_idx ) |
|
inline |
◆ SetSrc1SliceOrigin()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrc1SliceOrigin |
( |
const Src1Desc & | src1_desc, |
|
|
const Index & | src1_slice_origin_idx ) |
|
inline |
◆ SetSrc2SliceOrigin()
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| __device__ void ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::SetSrc2SliceOrigin |
( |
const Src2Desc & | src2_desc, |
|
|
const Index & | src2_slice_origin_idx ) |
|
inline |
◆ I0
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| auto ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::I0 = Number<0>{} |
|
staticconstexpr |
◆ nDim
template<typename Src0Data, typename Src1Data, typename Src2Data, typename DstData, typename Src0Desc, typename Src1Desc, typename Src2Desc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder,
index_t VectorDim,
index_t ScalarPerVector,
InMemoryDataOperationEnum DstInMemOp, bool Src0ResetCoordinateAfterRun, bool Src1ResetCoordinateAfterRun, bool Src2ResetCoordinateAfterRun, bool DstResetCoordinateAfterRun>
| index_t ck::ThreadwiseTensorSliceTransfer_v6r3< Src0Data, Src1Data, Src2Data, DstData, Src0Desc, Src1Desc, Src2Desc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, DstInMemOp, Src0ResetCoordinateAfterRun, Src1ResetCoordinateAfterRun, Src2ResetCoordinateAfterRun, DstResetCoordinateAfterRun >::nDim = SliceLengths::Size() |
|
staticconstexpr |
The documentation for this struct was generated from the following file: