BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack > Struct Template Reference#
Public Types |
Public Member Functions |
Static Public Member Functions |
Public Attributes |
Static Public Attributes |
Protected Types |
Protected Attributes |
Static Protected Attributes |
List of all members
ck::BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack > Struct Template Reference
#include <blockwise_gemm_dpp.hpp>
Public Types | |
| using | ThisThreadBlock = ThisThreadBlock<BlockSize> |
Public Member Functions | |
| __host__ __device__ constexpr auto & | GetCThreadBuffer () |
| __host__ __device__ | BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2 () |
| template<typename ABlockBuffer, typename BBlockBuffer, typename CThreadBuffer> | |
| __device__ void | Run (const ABlockBuffer &a_block_buf, const BBlockBuffer &b_block_buf, CThreadBuffer &c_thread_buf) const |
Static Public Member Functions | |
| static __device__ auto | GetWaveIdx () |
| static __device__ auto | CalculateAThreadOriginDataIndex_M0_M1_M2_K () |
| static __device__ auto | CalculateBThreadOriginDataIndex_N0_N1_N2_K () |
| template<index_t m0, index_t n0> | |
| static __device__ auto | CalculateCThreadOriginDataIndex (Number< m0 >, Number< n0 >) |
| __host__ static __device__ constexpr auto | GetCThreadDescriptor_M0_N0_M1_N1_M2_N2 () |
| __host__ static __device__ constexpr auto | GetCThreadDescriptor_G_M0_N0_M1_N1_M2_N2 () |
| __host__ static __device__ constexpr auto | GetCBlockDescriptor_M0_N0_M1_N1_M2_N2 () |
| __host__ static __device__ constexpr auto | GetCBlockDescriptor_G_M0_N0_M1_N1_M2_N2 () |
| template<typename CGridDesc_M_N> | |
| __host__ static __device__ constexpr auto | MakeCGridDescriptor_M0_N0_M1_N1_M2_N2 (const CGridDesc_M_N &c_grid_desc_m_n) |
| template<typename CGridDesc_G_M_N> | |
| __host__ static __device__ constexpr auto | MakeCGridDescriptor_G_M0_N0_M1_N1_M2_N2 (const CGridDesc_G_M_N &c_grid_desc_g_m_n) |
| __host__ static __device__ constexpr auto | MakeABlockDescriptor_M0_M1_M2_K () |
| __host__ static __device__ constexpr auto | MakeBBlockDescriptor_N0_N1_N2_K () |
Public Attributes | |
| StaticBufferTupleOfVector< AddressSpaceEnum::Vgpr, AccDataType, MRepeat *NRepeat, dpp_gemm.GetRegSizePerDpp(), true > | c_thread_buf_ |
Static Public Attributes | |
| static constexpr auto | I0 = Number<0>{} |
| static constexpr auto | I1 = Number<1>{} |
| static constexpr auto | I2 = Number<2>{} |
| static constexpr auto | I3 = Number<3>{} |
| static constexpr index_t | MPerBlock = AK0MK1BlockDesc{}.GetLength(I1) |
| static constexpr index_t | NPerBlock = BK0NK1BlockDesc{}.GetLength(I1) |
| static constexpr index_t | KPerBlock |
| static constexpr index_t | MWaves = MPerBlock / (MRepeat * MPerDpp) |
| static constexpr index_t | NWaves = NPerBlock / (NRepeat * NPerDpp) |
| static constexpr index_t | WaveSize = BlockSize / MWaves / NWaves |
| static constexpr index_t | A_K0 = AK0MK1BlockDesc{}.GetLength(I0) |
| static constexpr index_t | B_K0 = BK0NK1BlockDesc{}.GetLength(I0) |
| static constexpr index_t | A_K1 = AK0MK1BlockDesc{}.GetLength(I2) |
| static constexpr index_t | B_K1 = BK0NK1BlockDesc{}.GetLength(I2) |
| static constexpr auto | dpp_gemm = DppGemm<ABDataType, MPerDpp, NPerDpp, KPack>{} |
| static constexpr index_t | KPerThread = KPerBlock / dpp_gemm.K0PerDpp |
| static constexpr auto | a_block_desc_m0_m1_m2_k = MakeABlockDescriptor_M0_M1_M2_K() |
| static constexpr auto | b_block_desc_n0_n1_n2_k = MakeBBlockDescriptor_N0_N1_N2_K() |
Protected Types | |
| using | AThreadCopy |
| using | BThreadCopy |
Protected Attributes | |
| AThreadCopy | a_thread_copy_ {CalculateAThreadOriginDataIndex_M0_M1_M2_K()} |
| BThreadCopy | b_thread_copy_ {CalculateBThreadOriginDataIndex_N0_N1_N2_K()} |
Static Protected Attributes | |
| static constexpr auto | a_thread_desc_ |
| static constexpr auto | b_thread_desc_ |
| static constexpr auto | c_thread_desc_ |
Detailed Description
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
struct ck::BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack >
struct ck::BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack >
Blockwise GEMM that uses DPP instruction modifier to limit the amount of data loaded for each thread by sharing the data between threads in a lanegroup.
In every iteration, each wave calculates a C tile of size MPerDpp * NPerDpp, there are MRepeat iterations for M dimension and NRepeat for N one. In total, the algorithm runs using MPerBlock / (MRepeat * MPerDpp) * NPerBlock / (NRepeat * NPerDpp) waves.
Member Typedef Documentation
◆ AThreadCopy
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
protected |
Initial value:
ThreadwiseTensorSliceTransfer_v4<ABDataType,
ABDataType,
decltype(a_block_desc_m0_m1_m2_k),
decltype(a_thread_desc_),
3,
A_K1,
A_K1>
static constexpr auto a_block_desc_m0_m1_m2_k
Definition blockwise_gemm_dpp.hpp:254
static constexpr auto a_thread_desc_
Definition blockwise_gemm_dpp.hpp:312
static constexpr index_t A_K1
Definition blockwise_gemm_dpp.hpp:52
Definition utility/sequence.hpp:43
Definition threadwise_tensor_slice_transfer.hpp:1260
◆ BThreadCopy
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
protected |
Initial value:
ThreadwiseTensorSliceTransfer_v4<ABDataType,
ABDataType,
decltype(b_block_desc_n0_n1_n2_k),
decltype(b_thread_desc_),
3,
B_K1,
B_K1>
static constexpr index_t B_K1
Definition blockwise_gemm_dpp.hpp:53
static constexpr auto b_thread_desc_
Definition blockwise_gemm_dpp.hpp:316
static constexpr auto b_block_desc_n0_n1_n2_k
Definition blockwise_gemm_dpp.hpp:255
◆ ThisThreadBlock
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
| using ck::BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack >::ThisThreadBlock = ThisThreadBlock<BlockSize> |
Constructor & Destructor Documentation
◆ BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inline |
Member Function Documentation
◆ CalculateAThreadOriginDataIndex_M0_M1_M2_K()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestatic |
◆ CalculateBThreadOriginDataIndex_N0_N1_N2_K()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestatic |
◆ CalculateCThreadOriginDataIndex()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestatic |
◆ GetCBlockDescriptor_G_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ GetCBlockDescriptor_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ GetCThreadBuffer()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlineconstexpr |
◆ GetCThreadDescriptor_G_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ GetCThreadDescriptor_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ GetWaveIdx()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestatic |
◆ MakeABlockDescriptor_M0_M1_M2_K()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ MakeBBlockDescriptor_N0_N1_N2_K()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
inlinestaticconstexpr |
◆ MakeCGridDescriptor_G_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
template<typename CGridDesc_G_M_N>
|
inlinestaticconstexpr |
◆ MakeCGridDescriptor_M0_N0_M1_N1_M2_N2()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
template<typename CGridDesc_M_N>
|
inlinestaticconstexpr |
◆ Run()
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
template<typename ABlockBuffer, typename BBlockBuffer, typename CThreadBuffer>
|
inline |
Member Data Documentation
◆ a_block_desc_m0_m1_m2_k
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ A_K0
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ A_K1
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ a_thread_copy_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
protected |
◆ a_thread_desc_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexprprotected |
Initial value:
=
__host__ __device__ constexpr auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition tensor_descriptor_helper.hpp:101
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211
static constexpr auto I1
Definition blockwise_gemm_dpp.hpp:35
◆ b_block_desc_n0_n1_n2_k
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ B_K0
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ B_K1
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ b_thread_copy_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
protected |
◆ b_thread_desc_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexprprotected |
Initial value:
◆ c_thread_buf_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
| StaticBufferTupleOfVector<AddressSpaceEnum::Vgpr, AccDataType, MRepeat * NRepeat, dpp_gemm.GetRegSizePerDpp(), true> ck::BlockwiseGemmDpp_ak0mak1_bk0nbk1_m0n0m1n1m2n2< BlockSize, ABDataType, AccDataType, AK0MK1BlockDesc, BK0NK1BlockDesc, MPerDpp, NPerDpp, MRepeat, NRepeat, KPack >::c_thread_buf_ |
◆ c_thread_desc_
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexprprotected |
Initial value:
static constexpr auto dpp_gemm
Definition blockwise_gemm_dpp.hpp:55
◆ dpp_gemm
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ I0
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ I1
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ I2
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ I3
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ KPerBlock
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
Initial value:
=
static constexpr auto I2
Definition blockwise_gemm_dpp.hpp:36
static constexpr auto I0
Definition blockwise_gemm_dpp.hpp:34
◆ KPerThread
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ MPerBlock
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ MWaves
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ NPerBlock
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ NWaves
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
◆ WaveSize
template<index_t BlockSize, typename ABDataType, typename AccDataType, typename AK0MK1BlockDesc, typename BK0NK1BlockDesc, index_t MPerDpp, index_t NPerDpp, index_t MRepeat, index_t NRepeat, index_t KPack>
|
staticconstexpr |
The documentation for this struct was generated from the following file: