ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type > Struct Template Reference

ThreadwiseTensorSliceTransfer_v4&lt; SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type &gt; Struct Template Reference#

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

#include <threadwise_tensor_slice_transfer.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_v4 (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 SrcRefToOriginDisplacement, typename DstOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void Run (const SrcDesc &, const SrcRefToOriginDisplacement &, const SrcBuffer &src_buf, const DstData &scale, const DstDesc &, const DstOriginIdx &, DstBuffer &dst_buf) const
template<typename SrcSliceMoveStepIdx>
__device__ void MoveSrcSliceWindow (const SrcDesc &, const SrcSliceMoveStepIdx &src_slice_move_step_idx)
__device__ void SetSrcCoord (const Index &src_ref_idx)

Static Public Attributes

static constexpr index_t nDim = SliceLengths::Size()
static constexpr index_t PackedSize

Member Typedef Documentation

◆ Index

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::Index = MultiIndex<nDim>

◆ SrcCoord

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::SrcCoord = decltype(make_tensor_coordinate(SrcDesc{}, Index{}))

◆ SrcCoordStep

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::SrcCoordStep = decltype(make_tensor_coordinate_step(SrcDesc{}, Index{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v4()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::ThreadwiseTensorSliceTransfer_v4 ( const Index & src_ref_idx)
inlineconstexpr

Member Function Documentation

◆ MoveSrcSliceWindow()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcSliceMoveStepIdx>
__device__ void ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::MoveSrcSliceWindow ( const SrcDesc & ,
const SrcSliceMoveStepIdx & src_slice_move_step_idx )
inline

◆ Run() [1/2]

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcRefToOriginDisplacement, typename DstOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::Run ( const SrcDesc & ,
const SrcRefToOriginDisplacement & ,
const SrcBuffer & src_buf,
const DstData & scale,
const DstDesc & ,
const DstOriginIdx & ,
DstBuffer & dst_buf ) const
inline

◆ Run() [2/2]

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

◆ SetSrcCoord()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::SetSrcCoord ( const Index & src_ref_idx)
inline

Member Data Documentation

◆ nDim

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::nDim = SliceLengths::Size()
staticconstexpr

◆ PackedSize

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v4< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, type >::PackedSize
staticconstexpr
Initial value:
= []() {
if constexpr(is_same_v<remove_cvref_t<SrcData>, pk_i4_t>)
return 2;
else
return 1;
}()
ck::is_same_v
constexpr bool is_same_v
Definition type.hpp:283
ck::pk_i4_t
Definition data_type.hpp:187
The documentation for this struct was generated from the following file: