ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags > Struct Template Reference

ThreadwiseTensorSliceTransfer_v7&lt; SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags > Struct Template Reference

#include <threadwise_tensor_slice_transfer_v7.hpp>

Public Types

using Index = MultiIndex<nDim>
using SrcCoords = decltype(MakeCoordinates(SrcDescs{}, StaticallyIndexedArray<Index, nSrc>{}))
using DstCoords = decltype(MakeCoordinates(DstDescs{}, StaticallyIndexedArray<Index, nDst>{}))
using SpaceFillingCurve

Public Member Functions

__device__ constexpr ThreadwiseTensorSliceTransfer_v7 (const SrcDescs &src_descs, const StaticallyIndexedArray< Index, nSrc > &src_slice_origins, const DstDescs &dst_descs, const StaticallyIndexedArray< Index, nDst > &dst_slice_origins, const ElementwiseOperation &element_op)
template<typename Indices, enable_if_t< SrcDescs::Size()==Indices::Size(), bool > = false>
__device__ void SetSrcSliceOrigins (const SrcDescs &src_descs, const Indices &src_slice_origin_idxs)
template<typename Indices, enable_if_t< DstDescs::Size()==Indices::Size(), bool > = false>
__device__ void SetDstSliceOrigins (const DstDescs &dst_descs, const Indices &dst_slice_origin_idxs)
template<typename SrcBuffers, typename DstBuffers, enable_if_t< SrcDescs::Size()==SrcBuffers::Size() &&DstDescs::Size()==DstBuffers::Size(), bool > = false>
__device__ void Run (const SrcDescs &src_descs, const SrcBuffers &src_bufs, const DstDescs &dst_descs, DstBuffers dst_bufs)
template<index_t ISrc>
__device__ void MoveSrcSliceWindow (const SrcDescs &src_descs, Number< ISrc > iSrc, const Index &src_slice_origin_step_idx)
template<index_t IDst>
__device__ void MoveDstSliceWindow (const DstDescs &dst_descs, Number< IDst > iDst, const Index &dst_slice_origin_step_idx)

Static Public Member Functions

template<typename Descs, typename Indices, enable_if_t< Descs::Size()==Indices::Size(), bool > = false>
static constexpr auto MakeCoordinates (const Descs &descs, const Indices &indices)
static __device__ constexpr auto GetCoordinateResetStep ()

Static Public Attributes

static constexpr auto I0 = Number<0>{}
static constexpr index_t nDim = SliceLengths::Size()
static constexpr index_t nSrc = SrcDescs::Size()
static constexpr index_t nDst = DstDescs::Size()
static constexpr auto scalar_per_access

Member Typedef Documentation

◆ DstCoords

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
using ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::DstCoords = decltype(MakeCoordinates(DstDescs{}, StaticallyIndexedArray<Index, nDst>{}))

◆ Index

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
using ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::Index = MultiIndex<nDim>

◆ SpaceFillingCurve

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
using ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::SpaceFillingCurve
Initial value:
SpaceFillingCurve<SliceLengths, DimAccessOrder, remove_cv_t<decltype(scalar_per_access)>>
ck::remove_cv_t
typename remove_cv< T >::type remove_cv_t
Definition type.hpp:295
ck::ThreadwiseTensorSliceTransfer_v7::SpaceFillingCurve
SpaceFillingCurve< SliceLengths, DimAccessOrder, remove_cv_t< decltype(scalar_per_access)> > SpaceFillingCurve
Definition threadwise_tensor_slice_transfer_v7.hpp:69
ck::ThreadwiseTensorSliceTransfer_v7::scalar_per_access
static constexpr auto scalar_per_access
Definition threadwise_tensor_slice_transfer_v7.hpp:66

◆ SrcCoords

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
using ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::SrcCoords = decltype(MakeCoordinates(SrcDescs{}, StaticallyIndexedArray<Index, nSrc>{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v7()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::ThreadwiseTensorSliceTransfer_v7 ( const SrcDescs & src_descs,
const StaticallyIndexedArray< Index, nSrc > & src_slice_origins,
const DstDescs & dst_descs,
const StaticallyIndexedArray< Index, nDst > & dst_slice_origins,
const ElementwiseOperation & element_op )
inlineconstexpr

Member Function Documentation

◆ GetCoordinateResetStep()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::GetCoordinateResetStep ( )
inlinestaticconstexpr

◆ MakeCoordinates()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<typename Descs, typename Indices, enable_if_t< Descs::Size()==Indices::Size(), bool > = false>
constexpr auto ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::MakeCoordinates ( const Descs & descs,
const Indices & indices )
inlinestaticconstexpr

◆ MoveDstSliceWindow()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<index_t IDst>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::MoveDstSliceWindow ( const DstDescs & dst_descs,
Number< IDst > iDst,
const Index & dst_slice_origin_step_idx )
inline

◆ MoveSrcSliceWindow()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<index_t ISrc>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::MoveSrcSliceWindow ( const SrcDescs & src_descs,
Number< ISrc > iSrc,
const Index & src_slice_origin_step_idx )
inline

◆ Run()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<typename SrcBuffers, typename DstBuffers, enable_if_t< SrcDescs::Size()==SrcBuffers::Size() &&DstDescs::Size()==DstBuffers::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::Run ( const SrcDescs & src_descs,
const SrcBuffers & src_bufs,
const DstDescs & dst_descs,
DstBuffers dst_bufs )
inline

◆ SetDstSliceOrigins()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<typename Indices, enable_if_t< DstDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::SetDstSliceOrigins ( const DstDescs & dst_descs,
const Indices & dst_slice_origin_idxs )
inline

◆ SetSrcSliceOrigins()

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
template<typename Indices, enable_if_t< SrcDescs::Size()==Indices::Size(), bool > = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::SetSrcSliceOrigins ( const SrcDescs & src_descs,
const Indices & src_slice_origin_idxs )
inline

Member Data Documentation

◆ I0

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
auto ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::I0 = Number<0>{}
staticconstexpr

◆ nDim

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
index_t ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::nDim = SliceLengths::Size()
staticconstexpr

◆ nDst

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
index_t ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::nDst = DstDescs::Size()
staticconstexpr

◆ nSrc

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
index_t ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::nSrc = SrcDescs::Size()
staticconstexpr

◆ scalar_per_access

template<typename SrcDatas, typename DstDatas, typename SrcDescs, typename DstDescs, typename ElementwiseOperation, typename DstInMemOps, typename SliceLengths, typename DimAccessOrder, index_t VectorDim, index_t ScalarPerVector, typename SrcResetCoordinateAfterRunFlags, typename DstResetCoordinateAfterRunFlags>
auto ck::ThreadwiseTensorSliceTransfer_v7< SrcDatas, DstDatas, SrcDescs, DstDescs, ElementwiseOperation, DstInMemOps, SliceLengths, DimAccessOrder, VectorDim, ScalarPerVector, SrcResetCoordinateAfterRunFlags, DstResetCoordinateAfterRunFlags >::scalar_per_access
staticconstexpr
Initial value:
= generate_sequence(
detail::lambda_scalar_per_access<VectorDim, ScalarPerVector>{}, Number<nDim>{})
ck::Number
integral_constant< index_t, N > Number
Definition number.hpp:12
ck::generate_sequence
__host__ __device__ constexpr auto generate_sequence(F, Number< N >)
Definition sequence_helper.hpp:18
ck::detail::lambda_scalar_per_access
Definition threadwise_tensor_slice_transfer_util.hpp:20
The documentation for this struct was generated from the following file: