/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp Source File#
blockwise_gemm_pipeline_wmmaops_base.hpp
Go to the documentation of this file.
Definition: ck.hpp:267
typename detail::StaticallyIndexedArrayImpl< T, N >::type StaticallyIndexedArray
Definition: statically_indexed_array.hpp:45
__host__ constexpr __device__ auto make_multi_index(Xs &&... xs)
Definition: array_multi_index.hpp:15
__host__ constexpr __device__ auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition: tensor_descriptor_helper.hpp:49
__host__ constexpr __device__ auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition: tensor_descriptor_helper.hpp:101
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__host__ constexpr __device__ auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:425
__host__ constexpr __device__ auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:90
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:116
static constexpr index_t num_scale_krepeat
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:125
GridBuffer b_scale_grid_buf
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:162
__device__ BScale(GridDesc b_scale_grid_desc_, ThreadCopy b_scale_thread_copy_, GridBuffer b_scale_grid_buf_)
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:117
StaticallyIndexedArray< ThreadStaticBuffer, Number< NumberOfBuffers >{}> b_scale_thread_bufs
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:163
static constexpr auto b_scale_thread_copy_step
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:129
static constexpr index_t num_scale_k_block
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:124
static constexpr auto b_scale_thread_desc
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:127
__device__ void GlobalLoad(bool cond)
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:135
ThreadCopy b_scale_thread_copy
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:160
GridDesc b_scale_grid_desc
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:161
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:95
__device__ Empty()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:96
__device__ void GlobalLoad(bool cond)
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:98
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:35
AThreadCopy a_thread_copy_
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:377
__host__ constexpr __device__ auto & GetCThreadBuffer()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:166
static constexpr BWmmaTileDesc b_block_desc_k0_n0_n1_n2_k1
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:320
static constexpr auto I1
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:37
static constexpr index_t A_K1
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:57
static __device__ auto CalculateCThreadOriginDataIndex(Number< m0 >, Number< n0 >)
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:217
static __device__ auto CalculateBThreadOriginDataIndex()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:198
StaticBufferTupleOfVector< AddressSpaceEnum::Vgpr, AccDataType, MRepeat *NRepeat, wmma_gemm.GetRegSizePerWmma(), true > c_thread_buf_
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:92
__host__ __device__ BlockwiseGemmWmmaops_pipeline_base(Tuple6 a_origin=CalculateAThreadOriginDataIndex(), Tuple6 b_origin=CalculateBThreadOriginDataIndex())
Constructor for BlockwiseGemmWmmaops_pipeline_base.
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:264
static constexpr AWmmaTileDesc a_block_desc_k0_m0_m1_m2_k1
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:319
static constexpr index_t MWaves
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:46
static constexpr auto wmma_gemm
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:63
static constexpr index_t B_KRow
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:54
static __device__ auto GetWaveIdx()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:168
static constexpr auto I3
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:39
static constexpr auto I0
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:36
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:42
static constexpr index_t B_K1
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:58
__host__ static constexpr __device__ auto GetCThreadDescriptor_MRepeat_MWave_MSubGroup_NRepeat_NWave_NThreadPerSubGroup_MAccVgprs()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:281
static constexpr index_t WaveSize
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:44
__host__ static constexpr __device__ auto GetCBlockDescriptor_MRepeat_MWave_MSubGroup_NRepeat_NWave_NThreadPerSubGroup_MAccVgprs()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:302
static constexpr auto WmmaK
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:68
static constexpr auto I5
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:40
BThreadCopy b_thread_copy_
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:378
decltype(CalculateAThreadOriginDataIndex()) Tuple6
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:244
static constexpr index_t KRepeat
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:66
static constexpr index_t NWaves
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:47
static __device__ auto CalculateAThreadOriginDataIndex()
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:180
static constexpr index_t A_KRow
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:53
static constexpr auto I2
Definition: blockwise_gemm_pipeline_wmmaops_base.hpp:38
Definition: blockwise_gemm_pipeline_wmmaops.hpp:26
Definition: sequence.hpp:43
Definition: static_buffer.hpp:75
Definition: wmma_gemm.hpp:663
Definition: integral_constant.hpp:20
Definition: functional2.hpp:33