DppGemm< BaseType, MPerDpp, NPerDpp, KPack > Struct Template Reference

DppGemm&lt; BaseType, MPerDpp, NPerDpp, KPack &gt; Struct Template Reference#

Composable Kernel: ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack > Struct Template Reference

#include <dpp_gemm.hpp>

Public Types

using CIndex = MultiIndex<2>
using CIndex4D = MultiIndex<4>

Public Member Functions

__host__ __device__ constexpr DppGemm ()
template<class ADataType, class BDataType, class CDataType>
__device__ void Run (const ADataType &p_a_wave, const BDataType &p_b_wave, CDataType &p_c_thread) const

Static Public Member Functions

static __device__ constexpr index_t GetRegSizePerDpp ()
static __device__ auto GetLaneIdInWave ()
static __device__ auto GetWaveId ()
static __device__ auto GetLaneIdInLaneGroup ()
static __device__ auto GetLaneGroupIdInWave ()
static __device__ auto GetDppOpIdx ()
__host__ static __device__ auto CalculateAThreadOriginDataIndex_K_M ()
__host__ static __device__ auto CalculateBThreadOriginDataIndex_K_N ()
static __device__ CIndex GetBeginOfThreadBlk ()
__host__ static __device__ constexpr auto GetCMNThreadBlkLengths ()

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 auto I4 = Number<4>{}
static constexpr auto I5 = Number<5>{}
static constexpr auto dpp = DppSelector<BaseType, MPerDpp, NPerDpp>{}
static constexpr auto dpp_instr = dpp.selected_dpp
static constexpr auto K0PerDpp = 1
static constexpr auto K1PerDpp = dpp.GetK1PerDpp()

Member Typedef Documentation

◆ CIndex

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
using ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::CIndex = MultiIndex<2>

◆ CIndex4D

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
using ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::CIndex4D = MultiIndex<4>

Constructor & Destructor Documentation

◆ DppGemm()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__host__ __device__ constexpr ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::DppGemm ( )
inlineconstexpr

Member Function Documentation

◆ CalculateAThreadOriginDataIndex_K_M()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__host__ static __device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::CalculateAThreadOriginDataIndex_K_M ( )
inlinestatic

◆ CalculateBThreadOriginDataIndex_K_N()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__host__ static __device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::CalculateBThreadOriginDataIndex_K_N ( )
inlinestatic

◆ GetBeginOfThreadBlk()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ CIndex ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetBeginOfThreadBlk ( )
inlinestatic

◆ GetCMNThreadBlkLengths()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__host__ static __device__ constexpr auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetCMNThreadBlkLengths ( )
inlinestaticconstexpr

◆ GetDppOpIdx()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetDppOpIdx ( )
inlinestatic

◆ GetLaneGroupIdInWave()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetLaneGroupIdInWave ( )
inlinestatic

◆ GetLaneIdInLaneGroup()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetLaneIdInLaneGroup ( )
inlinestatic

◆ GetLaneIdInWave()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetLaneIdInWave ( )
inlinestatic

◆ GetRegSizePerDpp()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ constexpr index_t ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetRegSizePerDpp ( )
inlinestaticconstexpr

◆ GetWaveId()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
__device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetWaveId ( )
inlinestatic

◆ Run()

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
template<class ADataType, class BDataType, class CDataType>
__device__ void ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::Run ( const ADataType & p_a_wave,
const BDataType & p_b_wave,
CDataType & p_c_thread ) const
inline

Member Data Documentation

◆ dpp

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::dpp = DppSelector<BaseType, MPerDpp, NPerDpp>{}
staticconstexpr

◆ dpp_instr

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::dpp_instr = dpp.selected_dpp
staticconstexpr

◆ I0

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I2 = Number<2>{}
staticconstexpr

◆ I3

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I3 = Number<3>{}
staticconstexpr

◆ I4

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I4 = Number<4>{}
staticconstexpr

◆ I5

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I5 = Number<5>{}
staticconstexpr

◆ K0PerDpp

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::K0PerDpp = 1
staticconstexpr

◆ K1PerDpp

template<typename BaseType, index_t MPerDpp, index_t NPerDpp, index_t KPack>
auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::K1PerDpp = dpp.GetK1PerDpp()
staticconstexpr
The documentation for this struct was generated from the following file: