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__ constexpr __device__ 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 constexpr __device__ 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 constexpr __device__ 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__ constexpr __device__ 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>
static __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 constexpr __device__ auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::GetCMNThreadBlkLengths ( )
inlinestaticconstexpr

◆ GetDppOpIdx()

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

◆ GetLaneGroupIdInWave()

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

◆ GetLaneIdInLaneGroup()

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

◆ GetLaneIdInWave()

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

◆ GetRegSizePerDpp()

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

◆ GetWaveId()

template<typename BaseType , index_t MPerDpp, index_t NPerDpp, index_t KPack>
static __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>
constexpr 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>
constexpr 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>
constexpr auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::I0 = Number<0>{}
staticconstexpr

◆ I1

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

◆ I2

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

◆ I3

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

◆ I4

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

◆ I5

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

◆ K0PerDpp

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

◆ K1PerDpp

template<typename BaseType , index_t MPerDpp, index_t NPerDpp, index_t KPack>
constexpr auto ck::DppGemm< BaseType, MPerDpp, NPerDpp, KPack >::K1PerDpp = dpp.GetK1PerDpp()
staticconstexpr
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/tensor_operation/gpu/warp/dpp_gemm.hpp