/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/device/impl/split_k_offset_utils.hpp File Reference

• develop
• include
• ck
• tensor_operation
• gpu
• device
• impl

Classes | Namespaces | Functions

split_k_offset_utils.hpp File Reference

#include <numeric>
#include "ck/utility/common_header.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp"
#include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_selector.hpp"

Go to the source code of this file.

Classes
struct	ck::tensor_operation::device::SplitKHackEligibility< NDimSpatial, InLayout, WeiLayout, OutLayout >

Namespaces
	ck
	ck::tensor_operation
	ck::tensor_operation::device

Functions
template<typename Descriptor >
bool	ck::tensor_operation::device::IsDescriptorCompact (const Descriptor &desc)
template<typename GridwiseGemm , typename AGridDesc_AK0_M_K1 , typename BGridDesc_BK0_N_K1 , typename CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , bool HasMainKBlockLoop, InMemoryDataOperationEnum CGlobalMemoryDataOperation, TailNumber TailNum, typename ADataType , typename BDataType , typename CDataType >
__device__ void	ck::tensor_operation::device::DispatchSplitKHack (const ADataType *p_a_grid, const BDataType *p_b_grid, CDataType *p_c_grid, void *p_shared, const typename GridwiseGemm::Argument &karg, const AGridDesc_AK0_M_K1 &a_grid_desc_ak0_m_ak1, const BGridDesc_BK0_N_K1 &b_grid_desc_bk0_n_bk1, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &c_grid_desc_mblock_mperblock_nblock_nperblock, index_t k_id, index_t k_batch, bool split_k_offset_hack)
template<typename GridwiseGemm , typename AGridDesc_AK0_M_K1 , typename BGridDesc_BK0_N_K1 , typename CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , bool HasMainKBlockLoop, InMemoryDataOperationEnum CGlobalMemoryDataOperation, TailNumber TailNum, typename ADataType , typename BDataType , typename CDataType >
__device__ void	ck::tensor_operation::device::DispatchSplitKHack_2Lds (const ADataType *p_a_grid, const BDataType *p_b_grid, CDataType *p_c_grid, void *p_shared_0, void *p_shared_1, const typename GridwiseGemm::Argument &karg, const AGridDesc_AK0_M_K1 &a_grid_desc_ak0_m_ak1, const BGridDesc_BK0_N_K1 &b_grid_desc_bk0_n_bk1, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &c_grid_desc_mblock_mperblock_nblock_nperblock, index_t k_id, index_t k_batch, bool split_k_offset_hack)