ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type > Struct Template Reference

ThreadwiseTensorSliceTransfer_v2_gather&lt; SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type &gt; Struct Template Reference#

Composable Kernel: ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, 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_v2_gather (const SrcDesc &src_desc, const Index &src_slice_origin_idx, const StaticallyIndexedArray< index_t, scale_gather_num > &scale_gather_offsets)
__device__ void SetSrcSliceOrigin (const SrcDesc &src_desc, const Index &src_slice_origin_idx)
template<typename SrcBuffer, typename DstBuffer, typename DstSliceOriginIdx>
__device__ void Run (const SrcDesc &src_desc, const SrcBuffer &src_buf, const DstDesc &, const DstSliceOriginIdx &, DstBuffer &dst_buf)
__device__ void MoveSrcSliceWindow (const SrcDesc &src_desc, const Index &src_slice_origin_step_idx)
template<typename SrcMoveSliceWindowStepHack>
__device__ void MoveSrcSliceWindow (const SrcDesc &src_desc, const Index &src_slice_origin_step_idx, const SrcMoveSliceWindowStepHack &src_move_slice_window_step_hack)

Static Public Member Functions

static __device__ constexpr auto GetSrcCoordinateResetStep ()

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, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, 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, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, 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, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
using ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::SrcCoordStep = decltype(make_tensor_coordinate_step(SrcDesc{}, Index{}))

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_v2_gather()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::ThreadwiseTensorSliceTransfer_v2_gather ( const SrcDesc & src_desc,
const Index & src_slice_origin_idx,
const StaticallyIndexedArray< index_t, scale_gather_num > & scale_gather_offsets )
inlineconstexpr

Member Function Documentation

◆ GetSrcCoordinateResetStep()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ constexpr auto ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::GetSrcCoordinateResetStep ( )
inlinestaticconstexpr

◆ MoveSrcSliceWindow() [1/2]

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::MoveSrcSliceWindow ( const SrcDesc & src_desc,
const Index & src_slice_origin_step_idx )
inline

◆ MoveSrcSliceWindow() [2/2]

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcMoveSliceWindowStepHack>
__device__ void ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::MoveSrcSliceWindow ( const SrcDesc & src_desc,
const Index & src_slice_origin_step_idx,
const SrcMoveSliceWindowStepHack & src_move_slice_window_step_hack )
inline

◆ Run()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
template<typename SrcBuffer, typename DstBuffer, typename DstSliceOriginIdx>
__device__ void ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::Run ( const SrcDesc & src_desc,
const SrcBuffer & src_buf,
const DstDesc & ,
const DstSliceOriginIdx & ,
DstBuffer & dst_buf )
inline

◆ SetSrcSliceOrigin()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename SliceLengths, typename DimAccessOrder, index_t SrcVectorDim, index_t SrcScalarPerVector, index_t SrcScalarStrideInVector, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
__device__ void ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, type >::SetSrcSliceOrigin ( const SrcDesc & src_desc,
const Index & src_slice_origin_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, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, 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, bool SrcResetCoordinateAfterRun, index_t scale_gather_num, bool InvalidElementAsNaN = false, typename enable_if< DstDesc::IsKnownAtCompileTime(), bool >::type = false>
index_t ck::ThreadwiseTensorSliceTransfer_v2_gather< SrcData, DstData, SrcDesc, DstDesc, SliceLengths, DimAccessOrder, SrcVectorDim, SrcScalarPerVector, SrcScalarStrideInVector, SrcResetCoordinateAfterRun, scale_gather_num, InvalidElementAsNaN, 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: