/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp Source File#
cshuffle_epilogue.hpp
Go to the documentation of this file.
constexpr CK_TILE_HOST_DEVICE auto make_embed_tile_distribution_encoding(OuterDstr, InnerDstr)
Definition: tile_distribution_encoding.hpp:457
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_HOST_DEVICE auto make_naive_tensor_descriptor(const tuple< Lengths... > &lengths, const tuple< Strides... > &strides, number< GuaranteedLastDimensionVectorLength >=number<-1 >{}, number< GuaranteedLastDimensionVectorStride >=number<-1 >{})
Definition: tensor_descriptor.hpp:268
constexpr CK_TILE_HOST_DEVICE auto generate_tie(F &&f, number< N >)
Definition: tuple.hpp:435
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
CK_TILE_DEVICE auto tile_elementwise_inout_unpack(const InElementFunc &in_element_func, const Tuple &t, std::index_sequence< I... >)
Template function that "unpacks" a tuple and applies an element-wise operation.
Definition: tile_elementwise.hpp:71
constexpr CK_TILE_HOST_DEVICE auto to_sequence(tuple< number< Is >... >)
Definition: sequence.hpp:1052
constexpr CK_TILE_HOST_DEVICE auto merge_sequences(Seqs...)
Definition: sequence.hpp:823
constexpr CK_TILE_DEVICE auto make_tile_window(null_tensor_view, const WindowLengths &window_lengths, const multi_index< WindowLengths::size()> &, Ts &&...)
Definition: null_tile_window.hpp:75
CK_TILE_DEVICE void move_tile_window(null_tile_window< WindowLengths > &, const typename null_tile_window< WindowLengths >::BottomTensorIndex &)
Definition: null_tile_window.hpp:95
constexpr CK_TILE_HOST_DEVICE auto generate_tuple(F &&f, number< N >)
Definition: tuple.hpp:429
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
constexpr CK_TILE_HOST_DEVICE auto concat_tuple_of_reference(const tuple< X &... > &tx, const tuple< Y &... > &ty)
Definition: tuple.hpp:443
CK_TILE_DEVICE void update_tile(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: update_tile.hpp:22
CK_TILE_DEVICE void store_tile(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: store_tile.hpp:23
CK_TILE_DEVICE auto load_tile(const TileWindow_ &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={})
Definition: load_tile.hpp:22
constexpr CK_TILE_HOST_DEVICE auto make_static_tile_distribution(StaticTileDistributionEncoding_)
Definition: tile_distribution.hpp:480
typename impl::WarpGemmDispatcher< AType, BType, AccType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity, AttrNumAccess >::Type WarpGemmDispatcher
Definition: warp_gemm_dispatcher.hpp:178
typename uniform_sequence_gen< NSize, I >::type uniform_sequence_gen_t
Definition: sequence.hpp:1023
typename tuple_element< I, TTuple >::type tuple_element_t
Definition: tuple.hpp:208
Definition: cshuffle_epilogue.hpp:63
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:79
static constexpr CK_TILE_HOST_DEVICE auto MakeLdsBlockDescriptor()
Definition: cshuffle_epilogue.hpp:217
typename WG::CWarpTensor CWarpTensor
Definition: cshuffle_epilogue.hpp:214
remove_cvref_t< Problem_ > Problem
Definition: cshuffle_epilogue.hpp:64
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:84
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:88
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeD(number< I > index)
Get the vector store size for Di tensor.
Definition: cshuffle_epilogue.hpp:135
remove_cvref_t< typename Problem::ODataType > ODataType
Definition: cshuffle_epilogue.hpp:68
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:81
CK_TILE_DEVICE auto operator()(ODramWindow &out_dram_window, const OAccTile &o_acc_tile, const DsDramWindows &ds_dram_windows, void *p_smem)
Definition: cshuffle_epilogue.hpp:261
remove_cvref_t< typename Problem::ELayout > ELayout
Definition: cshuffle_epilogue.hpp:76
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:78
remove_cvref_t< typename Problem::DsLayout > DsLayout
Definition: cshuffle_epilogue.hpp:70
static constexpr CK_TILE_DEVICE auto MakeLdsDistributionEncode()
Definition: cshuffle_epilogue.hpp:239
static constexpr index_t MPerIteration
Definition: cshuffle_epilogue.hpp:90
static constexpr auto MNPerIterationShuffle
Definition: cshuffle_epilogue.hpp:194
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:87
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize()
Definition: cshuffle_epilogue.hpp:255
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeC()
Get the vector store size for C tensor.
Definition: cshuffle_epilogue.hpp:106
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:89
remove_cvref_t< typename Problem::DsDataType > DsDataType
Definition: cshuffle_epilogue.hpp:69
std::conditional_t< std::is_same_v< BDataType, pk_int4_t >, ADataType, BDataType > BTypeToUse
Definition: cshuffle_epilogue.hpp:75
remove_cvref_t< typename Problem::CDElementwise > CDElementwise
Definition: cshuffle_epilogue.hpp:77
static constexpr index_t NPerIterationShuffle
Definition: cshuffle_epilogue.hpp:203
remove_cvref_t< typename Problem::AccDataType > AccDataType
Definition: cshuffle_epilogue.hpp:67
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:92
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:86
static constexpr index_t NumMXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:191
static constexpr index_t NumNXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:192
WarpGemmDispatcher< ATypeToUse, BTypeToUse, AccDataType, MPerXdl, NPerXdl, KPerXdl, isCTransposed > WG
Definition: cshuffle_epilogue.hpp:211
static constexpr index_t MPerIterationShuffle
Definition: cshuffle_epilogue.hpp:202
remove_cvref_t< typename Problem::BDataType > BDataType
Definition: cshuffle_epilogue.hpp:66
static constexpr auto shuffle_tile_tuple
Shuffle tile configuration parameters.
Definition: cshuffle_epilogue.hpp:164
remove_cvref_t< typename Problem::ADataType > ADataType
Definition: cshuffle_epilogue.hpp:65
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:80
static constexpr index_t NPerIteration
Definition: cshuffle_epilogue.hpp:91
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:85
typename WG::CWarpDstr CWarpDstr
Definition: cshuffle_epilogue.hpp:213
std::conditional_t< std::is_same_v< ADataType, pk_int4_t >, BDataType, ADataType > ATypeToUse
Definition: cshuffle_epilogue.hpp:72
Definition: cshuffle_epilogue.hpp:33
remove_cvref_t< DsDataType_ > DsDataType
Definition: cshuffle_epilogue.hpp:38
remove_cvref_t< BDataType_ > BDataType
Definition: cshuffle_epilogue.hpp:35
remove_cvref_t< AccDataType_ > AccDataType
Definition: cshuffle_epilogue.hpp:36
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:48
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:42
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:43
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:49
remove_cvref_t< ELayout_ > ELayout
Definition: cshuffle_epilogue.hpp:40
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:47
remove_cvref_t< ADataType_ > ADataType
Definition: cshuffle_epilogue.hpp:34
static constexpr index_t MWave
Definition: cshuffle_epilogue.hpp:45
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:44
remove_cvref_t< CDElementwise_ > CDElementwise
Definition: cshuffle_epilogue.hpp:41
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:52
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:51
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:50
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:55
remove_cvref_t< DsLayout_ > DsLayout
Definition: cshuffle_epilogue.hpp:39
remove_cvref_t< ODataType_ > ODataType
Definition: cshuffle_epilogue.hpp:37
static constexpr index_t kNumWaveGroups
Definition: cshuffle_epilogue.hpp:54
static constexpr index_t NWave
Definition: cshuffle_epilogue.hpp:46
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:53
Definition: integral_constant.hpp:13
Definition: sequence.hpp:49
Definition: space_filling_curve.hpp:20
Definition: functional.hpp:43
Class creating 2D static tile distribution with different load/store patterns.
Definition: static_encoding_pattern.hpp:130
Definition: tile_distribution_encoding.hpp:26
Definition: tuple.hpp:192