/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.
405 // TODO: Check if there would be nicer ways to overload rather than with EmptyScale or nullptr_t
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:274
CK_TILE_DEVICE void tile_elementwise_inout(const InOutElementFunc &inout_element_func, InOutDstrTensors &... inout_dstr_tensors)
Definition: tile_elementwise.hpp:23
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:1055
constexpr CK_TILE_HOST_DEVICE auto merge_sequences(Seqs...)
Definition: sequence.hpp:826
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
typename detail::detector< nonesuch, void, Op, Args... >::value_t is_detected
Definition: type_traits.hpp:67
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:184
typename uniform_sequence_gen< NSize, I >::type uniform_sequence_gen_t
Definition: sequence.hpp:1026
typename tuple_element< I, TTuple >::type tuple_element_t
Definition: tuple.hpp:208
typename conditional< predicate, X, Y >::type conditional_t
Definition: functional.hpp:115
Definition: cshuffle_epilogue.hpp:407
typename T::DataType DataType
Definition: cshuffle_epilogue.hpp:419
Definition: cshuffle_epilogue.hpp:412
float DataType
Definition: cshuffle_epilogue.hpp:413
Definition: cshuffle_epilogue.hpp:68
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:99
CK_TILE_DEVICE void scale_tile(LdsTile &lds_tile, ScaleM &scale_m_window, ScaleN &scale_n_window)
Definition: cshuffle_epilogue.hpp:289
std::conditional_t< ADataTypeIsTuple, remove_cvref_t< AsDataType >, remove_cvref_t< tuple< AsDataType > >> AsDataTypeTuple
Definition: cshuffle_epilogue.hpp:82
static constexpr index_t NRepeat
Definition: cshuffle_epilogue.hpp:115
CK_TILE_DEVICE void slice_acc_tile(const OAccTile &o_acc_tile, LdsTile &lds_tile)
Definition: cshuffle_epilogue.hpp:330
static constexpr CK_TILE_HOST_DEVICE auto MakeLdsBlockDescriptor()
Definition: cshuffle_epilogue.hpp:244
static constexpr index_t MRepeat
Definition: cshuffle_epilogue.hpp:114
typename WG::CWarpTensor CWarpTensor
Definition: cshuffle_epilogue.hpp:237
typename WG::CWarpDstrEncoding CWarpDstrEncoding
Definition: cshuffle_epilogue.hpp:238
remove_cvref_t< typename Problem::AsDataType > AsDataType
Definition: cshuffle_epilogue.hpp:70
remove_cvref_t< Problem_ > Problem
Definition: cshuffle_epilogue.hpp:69
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:104
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:108
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeD(number< I > index)
Get the vector store size for Di tensor.
Definition: cshuffle_epilogue.hpp:158
remove_cvref_t< typename Problem::ODataType > ODataType
Definition: cshuffle_epilogue.hpp:73
CK_TILE_DEVICE void store_to_dram(OutDramWindow &out_dram_window, const COutTensor &c_out_tensor)
Definition: cshuffle_epilogue.hpp:371
static constexpr bool ADataTypeIsTuple
Definition: cshuffle_epilogue.hpp:77
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:101
remove_cvref_t< typename Problem::ELayout > ELayout
Definition: cshuffle_epilogue.hpp:96
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:98
static constexpr bool TiledMMAPermuteN
Definition: cshuffle_epilogue.hpp:109
static constexpr bool BDataTypeIsTuple
Definition: cshuffle_epilogue.hpp:78
remove_cvref_t< std::tuple_element_t< number< 0 >{}, BsDataTypeTuple > > BDataType
Definition: cshuffle_epilogue.hpp:89
remove_cvref_t< typename Problem::DsLayout > DsLayout
Definition: cshuffle_epilogue.hpp:75
static constexpr CK_TILE_DEVICE auto MakeLdsDistributionEncode()
Definition: cshuffle_epilogue.hpp:266
static constexpr index_t MPerIteration
Definition: cshuffle_epilogue.hpp:111
static constexpr auto MNPerIterationShuffle
Definition: cshuffle_epilogue.hpp:217
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:107
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize()
Definition: cshuffle_epilogue.hpp:282
CK_TILE_DEVICE void apply_d_tensors(DramWindows &d_dram_windows, COutTensor &c_out_tensor)
Definition: cshuffle_epilogue.hpp:357
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeC()
Get the vector store size for C tensor.
Definition: cshuffle_epilogue.hpp:129
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:110
remove_cvref_t< typename Problem::DsDataType > DsDataType
Definition: cshuffle_epilogue.hpp:74
std::conditional_t< std::is_same_v< BDataType, pk_int4_t >, ADataType, BDataType > BTypeToUse
Definition: cshuffle_epilogue.hpp:95
CK_TILE_DEVICE void move_windows(OutDramWindow &out_dram_window, DDramWindows &d_dram_windows)
Move both the output and D tensors windows for the next access.
Definition: cshuffle_epilogue.hpp:388
remove_cvref_t< typename Problem::CDElementwise > CDElementwise
Definition: cshuffle_epilogue.hpp:97
static constexpr index_t NPerIterationShuffle
Definition: cshuffle_epilogue.hpp:226
remove_cvref_t< typename Problem::AccDataType > AccDataType
Definition: cshuffle_epilogue.hpp:72
CK_TILE_DEVICE auto operator()(ODramWindow &out_dram_window, const OAccTile &o_acc_tile, const DsDramWindows &ds_dram_windows, void *, const ScaleM &scale_m={}, const ScaleN &scale_n={})
Definition: cshuffle_epilogue.hpp:429
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:113
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:106
CK_TILE_DEVICE auto operator()(ODramWindow &out_dram_window, const OAccTile &o_acc_tile, const DsDramWindows &ds_dram_windows, void *p_smem, const ScaleM &scale_m={}, const ScaleN &scale_n={})
Definition: cshuffle_epilogue.hpp:571
static constexpr index_t NumMXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:214
std::conditional_t< BDataTypeIsTuple, remove_cvref_t< BsDataType >, remove_cvref_t< tuple< BsDataType > >> BsDataTypeTuple
Definition: cshuffle_epilogue.hpp:86
static constexpr index_t NumNXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:215
remove_cvref_t< typename Problem::BsDataType > BsDataType
Definition: cshuffle_epilogue.hpp:71
WarpGemmDispatcher< ATypeToUse, BTypeToUse, AccDataType, MPerXdl, NPerXdl, KPerXdl, isCTransposed > WG
Definition: cshuffle_epilogue.hpp:234
static constexpr index_t MPerIterationShuffle
Definition: cshuffle_epilogue.hpp:225
CK_TILE_DEVICE void cast_lds_tile(LdsTile &lds_tile, InLdsWindow &in_lds_window)
Definition: cshuffle_epilogue.hpp:349
static constexpr auto shuffle_tile_tuple
Shuffle tile configuration parameters.
Definition: cshuffle_epilogue.hpp:187
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:100
static constexpr index_t NPerIteration
Definition: cshuffle_epilogue.hpp:112
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:105
typename WG::CWarpDstr CWarpDstr
Definition: cshuffle_epilogue.hpp:236
std::conditional_t< std::is_same_v< ADataType, pk_int4_t >, BDataType, ADataType > ATypeToUse
Definition: cshuffle_epilogue.hpp:92
remove_cvref_t< std::tuple_element_t< number< 0 >{}, AsDataTypeTuple > > ADataType
Definition: cshuffle_epilogue.hpp:88
Definition: cshuffle_epilogue.hpp:37
remove_cvref_t< AccDataType_ > AccDataType
Definition: cshuffle_epilogue.hpp:40
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:54
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:51
remove_cvref_t< CDElementwise_ > CDElementwise
Definition: cshuffle_epilogue.hpp:45
static constexpr bool TiledMMAPermuteN
Definition: cshuffle_epilogue.hpp:58
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:53
remove_cvref_t< AsDataType_ > AsDataType
Definition: cshuffle_epilogue.hpp:38
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:57
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:47
static constexpr index_t MWave
Definition: cshuffle_epilogue.hpp:49
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:56
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:60
remove_cvref_t< ODataType_ > ODataType
Definition: cshuffle_epilogue.hpp:41
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:52
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:48
remove_cvref_t< ELayout_ > ELayout
Definition: cshuffle_epilogue.hpp:44
remove_cvref_t< DsDataType_ > DsDataType
Definition: cshuffle_epilogue.hpp:42
static constexpr index_t NWave
Definition: cshuffle_epilogue.hpp:50
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:46
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:55
remove_cvref_t< DsLayout_ > DsLayout
Definition: cshuffle_epilogue.hpp:43
remove_cvref_t< BsDataType_ > BsDataType
Definition: cshuffle_epilogue.hpp:39
static constexpr index_t kNumWaveGroups
Definition: cshuffle_epilogue.hpp:59
Definition: integral_constant.hpp:13
Definition: unary_element_wise_operation.hpp:484
Definition: sequence.hpp:49
Definition: space_filling_curve.hpp:20
static constexpr CK_TILE_HOST_DEVICE auto get_forward_step(number< AccessIdx1d >)
Definition: space_filling_curve.hpp:70
static constexpr CK_TILE_HOST_DEVICE auto get_index(number< AccessIdx1d >)
Definition: space_filling_curve.hpp:158
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_access()
Definition: space_filling_curve.hpp:46
Definition: functional.hpp:43
Definition: tile_distribution_encoding.hpp:26
Definition: tuple.hpp:192