ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type > Struct Template Reference

ThreadwiseTensorSliceTransfer_v4r1&lt; SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type > Struct Template Reference
Public Types | Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v4r1.hpp>

Public Types

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

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v4r1 (const Index &src_ref_idx)
template<typename SrcRefToOriginDisplacement, typename DstOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void Run (const SrcDesc &, const SrcRefToOriginDisplacement &, const SrcBuffer &src_buf, const DstDesc &, const DstOriginIdx &, DstBuffer &dst_buf) const
template<typename SrcSliceMoveStepIdx>
__device__ void MoveSrcSliceWindow (const SrcDesc &, const SrcSliceMoveStepIdx &src_slice_move_step_idx)

Static Public Attributes

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

Member Typedef Documentation

◆ Index

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::Index = MultiIndex<nDim>

◆ SrcCoord

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::SrcCoord = decltype(make_tensor_coordinate(SrcDesc{}, Index{}))

◆ SrcCoordStep

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::SrcCoordStep = decltype(make_tensor_coordinate_step(SrcDesc{}, Index{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v4r1()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::ThreadwiseTensorSliceTransfer_v4r1 ( const Index & src_ref_idx)
inlineconstexpr

Member Function Documentation

◆ MoveSrcSliceWindow()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcSliceMoveStepIdx>
__device__ void ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::MoveSrcSliceWindow ( const SrcDesc & ,
const SrcSliceMoveStepIdx & src_slice_move_step_idx )
inline

◆ Run()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcRefToOriginDisplacement, typename DstOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::Run ( const SrcDesc & ,
const SrcRefToOriginDisplacement & ,
const SrcBuffer & src_buf,
const DstDesc & ,
const DstOriginIdx & ,
DstBuffer & dst_buf ) const
inline

Member Data Documentation

◆ I0

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
auto ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
auto ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::I1 = Number<1>{}
staticconstexpr

◆ nDim

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, typename SrcVectorTensorLengths, typename SrcVectorTensorContiguousDimOrder, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v4r1< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorTensorLengths, SrcVectorTensorContiguousDimOrder, type >::nDim = SliceLengths::Size()
staticconstexpr
The documentation for this struct was generated from the following file: