ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type > Struct Template Reference

• ck
• ThreadwiseTensorSliceTransfer_StaticToStatic

Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members

ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type > Struct Template Reference

Threadwise data transfer. More...

#include <threadwise_tensor_slice_transfer.hpp>

Public Types
using	Index = MultiIndex<nDim>

Public Member Functions
__device__ constexpr	ThreadwiseTensorSliceTransfer_StaticToStatic (const ElementwiseOperation &element_op)
template<typename SrcSliceOriginIdx, typename DstSliceOriginIdx, typename SrcBuffer, typename DstBuffer>
__device__ void	Run (const SrcDesc &, const SrcSliceOriginIdx &, const SrcBuffer &src_buf, const DstDesc &, const DstSliceOriginIdx &, DstBuffer &dst_buf) const

Public Attributes
ElementwiseOperation	element_op_

Static Public Attributes
static constexpr index_t	nDim = SliceLengths::Size()
static constexpr index_t	PackedSize

Detailed Description

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>
struct ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >

Threadwise data transfer.

Do NOT involve any tensor coordinates with StaticBuffer

Member Typedef Documentation

◆ Index

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

using ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >::Index = MultiIndex<nDim>

Constructor & Destructor Documentation

◆ ThreadwiseTensorSliceTransfer_StaticToStatic()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

__device__ constexpr ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >::ThreadwiseTensorSliceTransfer_StaticToStatic ( const ElementwiseOperation & element_op )

inlineconstexpr

Member Function Documentation

◆ Run()

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

template<typename SrcSliceOriginIdx, typename DstSliceOriginIdx, typename SrcBuffer, typename DstBuffer>

__device__ void ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >::Run ( const SrcDesc & , const SrcSliceOriginIdx & , const SrcBuffer & src_buf, const DstDesc & , const DstSliceOriginIdx & , DstBuffer & dst_buf ) const

inline

Member Data Documentation

◆ element_op_

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

ElementwiseOperation ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >::element_op_

◆ nDim

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

index_t ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, type >::nDim = SliceLengths::Size()

staticconstexpr

◆ PackedSize

template<typename SrcData, typename DstData, typename SrcDesc, typename DstDesc, typename ElementwiseOperation, typename SliceLengths, typename DimAccessOrder, index_t DstVectorDim, index_t DstScalarPerVector, typename enable_if< SrcDesc::IsKnownAtCompileTime() &&DstDesc::IsKnownAtCompileTime(), bool >::type = false>

index_t ck::ThreadwiseTensorSliceTransfer_StaticToStatic< SrcData, DstData, SrcDesc, DstDesc, ElementwiseOperation, SliceLengths, DimAccessOrder, DstVectorDim, DstScalarPerVector, 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:

• threadwise_tensor_slice_transfer.hpp