GridwiseGemmPipeline_v1_WeightOnly< 1, true, true > Struct Reference

GridwiseGemmPipeline_v1_WeightOnly< 1, true, true > Struct Reference#

Composable Kernel: ck::GridwiseGemmPipeline_v1_WeightOnly< 1, true, true > Struct Reference

#include <gridwise_gemm_pipeline_v1.hpp>

Static Public Member Functions
__host__ static __device__ constexpr bool	IsSupported (index_t)
__host__ static __device__ constexpr bool	CalculateHasMainLoop (index_t num_loop)
template<bool HasMainLoop, typename AGridDesc, typename ABlockDesc, typename ABlockTransfer, typename AGridBuffer, typename ABlockBuffer, typename ABlockTransferStep, typename BGridDesc, typename BBlockDesc, typename BBlockTransfer, typename BGridBuffer, typename BBlockBuffer, typename BBlockTransferStep, typename ScaleGridDesc, typename ScaleGridBuffer, typename BlockwiseGemm, typename CThreadBuffer>
static __device__ void	Run (const AGridDesc &a_grid_desc, const ABlockDesc &a_block_desc, ABlockTransfer &a_blockwise_copy, const AGridBuffer &a_grid_buf, ABlockBuffer &a_block_buf, const ABlockTransferStep &a_block_copy_step, const BGridDesc &b_grid_desc, const BBlockDesc &b_block_desc, BBlockTransfer &b_blockwise_copy, const BGridBuffer &b_grid_buf, BBlockBuffer &b_block_buf, const BBlockTransferStep &b_block_copy_step, const ScaleGridDesc &scale_grid_desc, const ScaleGridBuffer &scale_grid_buf, const BlockwiseGemm &blockwise_gemm, CThreadBuffer &c_thread_buf, index_t num_loop)

Static Public Attributes
static constexpr auto	I0 = Number<0>{}
static constexpr auto	I1 = Number<1>{}

Member Function Documentation

◆ CalculateHasMainLoop()

__host__ static __device__ constexpr bool ck::GridwiseGemmPipeline_v1_WeightOnly< 1, true, true >::CalculateHasMainLoop ( index_t num_loop )

inlinestaticconstexpr

◆ IsSupported()

__host__ static __device__ constexpr bool ck::GridwiseGemmPipeline_v1_WeightOnly< 1, true, true >::IsSupported ( index_t )

inlinestaticconstexpr

◆ Run()

template<bool HasMainLoop, typename AGridDesc, typename ABlockDesc, typename ABlockTransfer, typename AGridBuffer, typename ABlockBuffer, typename ABlockTransferStep, typename BGridDesc, typename BBlockDesc, typename BBlockTransfer, typename BGridBuffer, typename BBlockBuffer, typename BBlockTransferStep, typename ScaleGridDesc, typename ScaleGridBuffer, typename BlockwiseGemm, typename CThreadBuffer>

__device__ void ck::GridwiseGemmPipeline_v1_WeightOnly< 1, true, true >::Run	(	const AGridDesc &	a_grid_desc,
		const ABlockDesc &	a_block_desc,
		ABlockTransfer &	a_blockwise_copy,
		const AGridBuffer &	a_grid_buf,
		ABlockBuffer &	a_block_buf,
		const ABlockTransferStep &	a_block_copy_step,
		const BGridDesc &	b_grid_desc,
		const BBlockDesc &	b_block_desc,
		BBlockTransfer &	b_blockwise_copy,
		const BGridBuffer &	b_grid_buf,
		BBlockBuffer &	b_block_buf,
		const BBlockTransferStep &	b_block_copy_step,
		const ScaleGridDesc &	scale_grid_desc,
		const ScaleGridBuffer &	scale_grid_buf,
		const BlockwiseGemm &	blockwise_gemm,
		CThreadBuffer &	c_thread_buf,
		index_t	num_loop )