21 template <
typename GridwiseGemm,
27 typename CGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
28 typename AElementwiseOperation,
29 typename BElementwiseOperation,
30 typename CElementwiseOperation,
31 typename Block2CTileMap,
32 bool HasMainKBlockLoop>
34 #if CK_USE_LAUNCH_BOUNDS
38 const BDataType* __restrict__ p_b_grid,
39 CDataType* __restrict__ p_c_grid,
40 const AGridDesc a_grid_desc,
41 const BGridDesc b_grid_desc,
42 const CGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
43 c_grid_desc_mblock_mperblock_nblock_nperblock,
44 const AElementwiseOperation a_element_op,
45 const BElementwiseOperation b_element_op,
46 const CElementwiseOperation c_element_op,
47 const Block2CTileMap block_2_ctile_map)
49 #if(defined(__gfx11__) || defined(__gfx12__))
50 __shared__
char p_shared[GridwiseGemm::SharedMemTrait::lds_size];
52 GridwiseGemm::template Run<HasMainKBlockLoop>(p_a_grid,
58 c_grid_desc_mblock_mperblock_nblock_nperblock,
69 ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
73 ignore = block_2_ctile_map;
81 typename CShuffleDataType,
86 typename CGridDesc_M_N,
87 typename AElementwiseOperation,
88 typename BElementwiseOperation,
89 typename CElementwiseOperation,
98 typename ABlockTransferThreadClusterLengths_K0_M_K1,
99 typename ABlockTransferThreadClusterArrangeOrder,
100 typename ABlockTransferSrcAccessOrder,
101 index_t ABlockTransferSrcVectorDim,
102 index_t ABlockTransferSrcScalarPerVector,
103 index_t ABlockTransferDstScalarPerVector_K1,
104 bool AThreadTransferSrcResetCoordinateAfterRun,
106 bool ABlockLdsExtraM,
107 typename BBlockTransferThreadClusterLengths_K0_N_K1,
108 typename BBlockTransferThreadClusterArrangeOrder,
109 typename BBlockTransferSrcAccessOrder,
110 index_t BBlockTransferSrcVectorDim,
111 index_t BBlockTransferSrcScalarPerVector,
112 index_t BBlockTransferDstScalarPerVector_K1,
113 bool BThreadTransferSrcResetCoordinateAfterRun,
115 bool BBlockLdsExtraN,
116 index_t CShuffleMRepeatPerShuffle,
117 index_t CShuffleNRepeatPerShuffle,
118 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
119 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
120 index_t NumGemmKPrefetchStage = 1,
137 static constexpr
auto MWaves = MPerBlock / (MRepeat * MPerWmma);
138 static constexpr
auto NWaves = NPerBlock / (NRepeat * NPerWmma);
139 static constexpr
auto WmmaK =
K1 == 16 ? 32 : 16;
145 NumGemmKPrefetchStage,
153 constexpr
auto a_block_desc = [&]() {
154 if constexpr(AEnableLds)
157 constexpr
auto K0PerBlock = KPerBlock /
K1;
158 constexpr
auto max_lds_align =
K1;
160 if constexpr(ABlockLdsExtraM)
174 constexpr
auto A_KRow =
I2;
175 constexpr
auto KWmmaPerblock = KPerBlock /
WmmaK;
176 constexpr
auto K0PerWmma =
WmmaK / A_KRow /
K1;
201 constexpr
auto b_block_desc = [&]() {
202 if constexpr(BEnableLds)
205 constexpr
auto K0PerBlock = KPerBlock /
K1;
206 constexpr
auto max_lds_align =
K1;
208 if constexpr(BBlockLdsExtraN)
223 constexpr
auto B_KRow =
I2;
224 constexpr
auto KWmmaPerblock = KPerBlock /
WmmaK;
225 constexpr
auto K0PerWmma =
WmmaK / B_KRow /
K1;
250 constexpr
auto a_block_copy_step = [&]() {
251 if constexpr(AEnableLds)
253 constexpr
auto K0PerBlock = KPerBlock /
K1;
259 constexpr
auto KWmmaPerBlock = KPerBlock /
WmmaK;
265 return a_block_copy_step;
270 constexpr
auto b_block_copy_step = [&]() {
271 if constexpr(BEnableLds)
273 constexpr
auto K0PerBlock = KPerBlock /
K1;
279 constexpr
auto KWmmaPerBlock = KPerBlock /
WmmaK;
285 return b_block_copy_step;
289 template <
typename ABlockDesc_>
293 constexpr
auto a_wave_desc = [&]() {
294 if constexpr(AEnableLds)
297 constexpr
auto A_K0 = ABlockDesc_{}.GetLength(
I0);
298 constexpr
auto A_K1 = ABlockDesc_{}.GetLength(
I2);
300 constexpr
auto A_KRow =
I2;
302 constexpr
auto A_KRow =
I1;
317 constexpr
auto KWmma = ABlockDesc_{}.GetLength(
I0);
318 constexpr
auto K0PerWmma = ABlockDesc_{}.GetLength(
I3);
319 constexpr
auto A_KRow = ABlockDesc_{}.GetLength(
I4);
320 constexpr
auto A_K1 = ABlockDesc_{}.GetLength(
I6);
353 template <
typename BBlockDesc_>
356 constexpr
auto b_wave_desc = [&]() {
357 if constexpr(BEnableLds)
360 constexpr
auto B_K0 = BBlockDesc_{}.GetLength(
I0);
361 constexpr
auto B_K1 = BBlockDesc_{}.GetLength(
I2);
363 constexpr
auto B_KRow =
I2;
365 constexpr
auto B_KRow =
I1;
379 constexpr
auto KWmma = BBlockDesc_{}.GetLength(
I0);
380 constexpr
auto K0PerWmma = BBlockDesc_{}.GetLength(
I3);
381 constexpr
auto B_KRow = BBlockDesc_{}.GetLength(
I4);
382 constexpr
auto B_K1 = BBlockDesc_{}.GetLength(
I6);
397 __host__ __device__
static constexpr
auto
401 constexpr
auto c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat =
408 return c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat;
412 template <
typename Block2CTileMap>
413 __host__ __device__
static constexpr
bool CheckValidity(
const AGridDesc& a_grid_desc,
414 const BGridDesc& b_grid_desc,
415 const CGridDesc_M_N& c_grid_desc_m_n,
416 const Block2CTileMap& block_2_ctile_map)
419 "wrong! K1 need to be known at compile-time");
421 static_assert((MPerBlock % (MPerWmma * MRepeat) == 0) &&
422 (NPerBlock % (NRepeat * NPerWmma)) == 0,
423 "Invalid tuning param!");
425 const auto GetAProblemsizeMK = [&]() {
426 if constexpr(AEnableLds)
429 a_grid_desc.GetLength(
I0) * a_grid_desc.GetLength(
I2));
433 return make_tuple(a_grid_desc.GetLength(
I1) * a_grid_desc.GetLength(
I2) *
434 a_grid_desc.GetLength(
I5),
435 a_grid_desc.GetLength(
I0) * a_grid_desc.GetLength(
I3) *
436 a_grid_desc.GetLength(
I4) * a_grid_desc.GetLength(
I6));
440 const auto GetBProblemsizeNK = [&]() {
441 if constexpr(BEnableLds)
444 b_grid_desc.GetLength(
I0) * b_grid_desc.GetLength(
I2));
448 return make_tuple(b_grid_desc.GetLength(
I1) * b_grid_desc.GetLength(
I2) *
449 b_grid_desc.GetLength(
I5),
450 b_grid_desc.GetLength(
I0) * b_grid_desc.GetLength(
I3) *
451 b_grid_desc.GetLength(
I4) * b_grid_desc.GetLength(
I6));
455 const auto M = GetAProblemsizeMK()[
I0];
456 const auto N = GetBProblemsizeNK()[
I0];
457 const auto K = GetAProblemsizeMK()[
I1];
459 if(!(M == c_grid_desc_m_n.GetLength(
I0) && N == c_grid_desc_m_n.GetLength(
I1) &&
460 K == GetBProblemsizeNK()[
I1]))
464 printf(
"A: MxK = %d x %d, B: NxK = %d x %d, C: MxN = %d x %d\n",
465 GetAProblemsizeMK()[
I0],
466 GetAProblemsizeMK()[
I1],
467 GetBProblemsizeNK()[
I0],
468 GetBProblemsizeNK()[
I1],
469 c_grid_desc_m_n.GetLength(
I0),
470 c_grid_desc_m_n.GetLength(
I1));
471 printf(
"GridwiseOp err: ProblemSize check");
476 if(!(M % MPerBlock == 0 && N % NPerBlock == 0 && K % KPerBlock == 0))
480 printf(
"GridwiseOp err: ProblemSize division");
486 const auto num_k_loop = K / KPerBlock;
488 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
492 printf(
"GridwiseOp err: Pipeline not support this k_loop");
497 if(!block_2_ctile_map.CheckValidity(c_grid_desc_m_n))
505 if(!(a_grid_desc.GetElementSpaceSize() *
sizeof(ADataType) <= TwoGB &&
506 b_grid_desc.GetElementSpaceSize() *
sizeof(BDataType) <= TwoGB))
515 const index_t num_loop = K / KPerBlock;
517 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
520 __host__ __device__
static constexpr
auto
523 const auto M = c_grid_desc_m_n.GetLength(
I0);
524 const auto N = c_grid_desc_m_n.GetLength(
I1);
526 const auto MBlock = M / MPerBlock;
527 const auto NBlock = N / NPerBlock;
536 return c_grid_desc_mblock_mperblock_nblock_nperblock;
568 .GetElementSpaceSize();
584 template <
bool HasMainKBlockLoop,
typename Block2CTileMap = DefaultBlock2CTileMap>
585 __device__
static void Run(
const ADataType* __restrict__ p_a_grid,
586 const BDataType* __restrict__ p_b_grid,
587 CDataType* __restrict__ p_c_grid,
588 void* __restrict__ p_shared,
589 const AGridDesc& a_grid_desc,
590 const BGridDesc& b_grid_desc,
592 c_grid_desc_mblock_mperblock_nblock_nperblock,
593 const AElementwiseOperation& a_element_op,
594 const BElementwiseOperation& b_element_op,
595 const CElementwiseOperation& c_element_op,
596 const Block2CTileMap& block_2_ctile_map)
601 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
602 p_a_grid, a_grid_desc.GetElementSpaceSize());
603 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
604 p_b_grid, b_grid_desc.GetElementSpaceSize());
605 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
606 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
611 if(!block_2_ctile_map.ValidCTileIndex(
613 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
614 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
618 const index_t m_block_data_idx_on_grid = __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
619 const index_t n_block_data_idx_on_grid = __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
623 const auto K = [&](){
624 if constexpr(AEnableLds){
625 return a_grid_desc.GetLength(
I0) * a_grid_desc.GetLength(
I2);
628 return a_grid_desc.GetLength(
I0) * a_grid_desc.GetLength(
I3)
629 * a_grid_desc.GetLength(
I4) * a_grid_desc.GetLength(
I6);
636 auto a_block_trait = [&](){
638 if constexpr(AEnableLds)
640 constexpr
auto K0PerBlock = KPerBlock/
K1;
641 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
642 static_cast<ADataType*
>(p_shared),
645 auto a_blockwise_copy =
647 AElementwiseOperation,
651 ABlockTransferThreadClusterLengths_K0_M_K1,
652 ABlockTransferThreadClusterArrangeOrder,
655 decltype(a_grid_desc),
656 decltype(a_block_desc),
657 ABlockTransferSrcAccessOrder,
659 ABlockTransferSrcVectorDim,
661 ABlockTransferSrcScalarPerVector,
662 ABlockTransferDstScalarPerVector_K1,
665 AThreadTransferSrcResetCoordinateAfterRun,
667 NumGemmKPrefetchStage>(
675 return make_tuple(a_block_buf, a_blockwise_copy);
681 constexpr
auto KWmmaPerBlock = KPerBlock /
WmmaK;
682 constexpr
auto K0PerWmma =
WmmaK/2/K1Value;
683 auto a_block_buf = make_static_buffer<AddressSpaceEnum::Vgpr, ADataType>(
684 a_block_desc.GetElementSpaceSize());
687 auto a_blockwise_copy =
690 decltype(a_grid_desc),
691 decltype(a_block_desc),
701 ABlockTransferSrcScalarPerVector,
702 AThreadTransferSrcResetCoordinateAfterRun,
706 m_block_data_idx_on_grid/(
MWaves * MPerWmma),
713 return make_tuple(a_block_buf, a_blockwise_copy);
717 auto b_block_trait = [&](){
718 if constexpr(BEnableLds)
720 constexpr
auto K0PerBlock = KPerBlock/
K1;
721 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
725 auto b_blockwise_copy =
727 BElementwiseOperation,
731 BBlockTransferThreadClusterLengths_K0_N_K1,
732 BBlockTransferThreadClusterArrangeOrder,
735 decltype(b_grid_desc),
736 decltype(b_block_desc),
737 BBlockTransferSrcAccessOrder,
739 BBlockTransferSrcVectorDim,
741 BBlockTransferSrcScalarPerVector,
742 BBlockTransferDstScalarPerVector_K1,
745 BThreadTransferSrcResetCoordinateAfterRun,
747 NumGemmKPrefetchStage>(
755 return make_tuple(b_block_buf, b_blockwise_copy);
761 constexpr
auto KWmmaPerBlock = KPerBlock /
WmmaK;
762 constexpr
auto K0PerWmma =
WmmaK/2/K1Value;
763 auto b_block_buf = make_static_buffer<AddressSpaceEnum::Vgpr, BDataType>(
764 b_block_desc.GetElementSpaceSize());
767 auto b_blockwise_copy =
770 decltype(b_grid_desc),
771 decltype(b_block_desc),
781 BBlockTransferSrcScalarPerVector,
782 BThreadTransferSrcResetCoordinateAfterRun,
786 n_block_data_idx_on_grid/(
NWaves * NPerWmma),
793 return make_tuple(b_block_buf, b_blockwise_copy);
797 auto a_block_buf = a_block_trait()[
I0];
798 auto a_blockwise_copy = a_block_trait()[
I1];
800 auto b_block_buf = b_block_trait()[
I0];
801 auto b_blockwise_copy = b_block_trait()[
I1];
806 auto blockwise_gemm =
833 const index_t KBlockMainLoop = __builtin_amdgcn_readfirstlane(K / KPerBlock);
834 GridwiseGemmPipe::template Run<HasMainKBlockLoop>(a_grid_desc,
839 a_block_slice_copy_step,
845 b_block_slice_copy_step,
853 constexpr
auto c_thread_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs =
854 blockwise_gemm.GetCThreadDescriptor_MRepeat_MWave_MSubGroup_NRepeat_NWave_NThreadPerSubGroup_MAccVgprs();
857 constexpr
auto c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp =
858 blockwise_gemm.GetCBlockDescriptor_MRepeat_MWave_MSubGroup_NRepeat_NWave_NThreadPerSubGroup_MAccVgprs();
860 constexpr
auto MWave = c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp.GetLength(
I1);
861 constexpr
auto MSubGroup = c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp.GetLength(
I2);
862 constexpr
auto NWave = c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp.GetLength(
I4);
863 constexpr
auto NThreadPerSubGroup = c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp.GetLength(
I5);
864 constexpr
auto MAccVgprs = c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs_tmp.GetLength(
I6);
867 constexpr
auto c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat =
870 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
875 c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat,
887 NThreadPerSubGroup))),
893 const auto c_thread_mtx_on_block = blockwise_gemm.CalculateCThreadOriginDataIndex(
I0,
I0);
895 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
896 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
898 const auto m_thread_data_on_block_to_mrepeat_mwave_msubgroup_maccvgprs_adaptor =
904 const auto n_thread_data_on_block_to_nrepeat_nwave_nthreadpersubgroup_adaptor =
910 const auto m_thread_data_on_block_idx = m_thread_data_on_block_to_mrepeat_mwave_msubgroup_maccvgprs_adaptor.CalculateBottomIndex(
913 const auto n_thread_data_on_block_idx = n_thread_data_on_block_to_nrepeat_nwave_nthreadpersubgroup_adaptor.CalculateBottomIndex(
917 auto c_thread_copy_vgpr_to_lds =
920 decltype(c_thread_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs),
921 decltype(c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs),
926 CShuffleNRepeatPerShuffle,
936 c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs,
938 m_thread_data_on_block_idx[
I1],
939 m_thread_data_on_block_idx[
I2],
941 n_thread_data_on_block_idx[
I1],
942 n_thread_data_on_block_idx[
I2],
943 m_thread_data_on_block_idx[
I3]),
949 CElementwiseOperation,
950 CGlobalMemoryDataOperation,
952 CShuffleMRepeatPerShuffle * MWave * MPerWmma,
954 CShuffleNRepeatPerShuffle * NWave * NPerWmma>,
955 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
959 decltype(c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat),
960 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
963 CShuffleBlockTransferScalarPerVector_NPerBlock,
966 {c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat,
968 c_grid_desc_mblock_mperblock_nblock_nperblock,
974 constexpr
auto sfc_c_vgpr =
980 CShuffleNRepeatPerShuffle,
986 constexpr
auto sfc_c_global =
990 CShuffleMRepeatPerShuffle * MWave * MPerWmma,
992 CShuffleNRepeatPerShuffle * NWave * NPerWmma>>{};
994 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
996 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1003 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs,
1004 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1006 c_block_desc_mrepeat_mwave_msubgroup_nrepeat_nwave_nthreadpersubgroup_maccvgprs,
1007 c_shuffle_block_buf);
1013 c_shuffle_block_copy_lds_to_global.Run(
1014 c_shuffle_block_desc_mshrepeat_mpershrepeat_nshrepeat_npershrepeat,
1015 c_shuffle_block_buf,
1016 c_grid_desc_mblock_mperblock_nblock_nperblock,
1019 if constexpr(access_id < num_access - 1)
1021 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1024 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1025 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:31
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:30
__host__ constexpr __device__ auto integer_least_multiple(X x, Y y)
Definition: math.hpp:78
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
__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
constexpr auto GridwiseGemmPipeline_Selector()
Definition: gridwise_gemm_pipeline_selector.hpp:31
InMemoryDataOperationEnum
Definition: ck.hpp:277
__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
__global__ void kernel_gemm_wmma(const ADataType *__restrict__ p_a_grid, const BDataType *__restrict__ p_b_grid, CDataType *__restrict__ p_c_grid, const AGridDesc a_grid_desc, const BGridDesc b_grid_desc, const CGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock, const AElementwiseOperation a_element_op, const BElementwiseOperation b_element_op, const CElementwiseOperation c_element_op, const Block2CTileMap block_2_ctile_map)
Definition: gridwise_gemm_wmma.hpp:37
int64_t long_index_t
Definition: ck.hpp:300
__host__ constexpr __device__ auto make_naive_tensor_descriptor_aligned(const Tuple< Lengths... > &lengths, Align align)
Definition: tensor_descriptor_helper.hpp:132
__host__ constexpr __device__ auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:425
__host__ constexpr __device__ auto make_freeze_transform(const LowerIndex &low_idx)
Definition: multi_index_transform_helper.hpp:151
constexpr detail::ignore_t ignore
Definition: ignore.hpp:20
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:47
bool EnvIsEnabled(EnvVar)
Definition: env.hpp:140
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__host__ constexpr __device__ auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:211
remove_cv_t< remove_reference_t< T > > remove_cvref_t
Definition: type.hpp:297
__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
LoopScheduler
Definition: loop_scheduler.hpp:15
int32_t index_t
Definition: ck.hpp:299
__device__ index_t get_thread_local_1d_id()
Definition: get_id.hpp:41
__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_descriptor.hpp:319
__device__ void block_sync_lds()
Definition: synchronization.hpp:10
PipelineVersion
Definition: gridwise_gemm_pipeline_selector.hpp:18
typename remove_cv< T >::type remove_cv_t
Definition: type.hpp:295
constexpr LoopScheduler make_default_loop_scheduler()
Definition: loop_scheduler.hpp:20
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
Definition: block_to_ctile_map.hpp:261
Definition: blockwise_gemm_wmma.hpp:550
__host__ constexpr __device__ auto & GetCThreadBuffer()
Definition: blockwise_gemm_wmma.hpp:585
Definition: gridwise_gemm_wmma.hpp:548
static constexpr auto c_shuffle_block_space_size
Definition: gridwise_gemm_wmma.hpp:566
static constexpr auto b_block_space_size_aligned
Definition: gridwise_gemm_wmma.hpp:557
static constexpr auto max_lds_align
Definition: gridwise_gemm_wmma.hpp:551
static constexpr auto c_shuffle_block_space_offset
Definition: gridwise_gemm_wmma.hpp:570
static constexpr auto lds_size
Definition: gridwise_gemm_wmma.hpp:572
static constexpr auto a_block_space_size_aligned
Definition: gridwise_gemm_wmma.hpp:553
static constexpr auto a_block_space_offset
Definition: gridwise_gemm_wmma.hpp:562
static constexpr auto b_block_space_offset
Definition: gridwise_gemm_wmma.hpp:563
Definition: gridwise_gemm_wmma.hpp:124
__host__ static constexpr __device__ auto MakeDefaultBlock2CTileMap(const CGridDesc_M_N &c_grid_desc_m_n, index_t, index_t)
Definition: gridwise_gemm_wmma.hpp:540
remove_cvref_t< decltype(MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(CGridDesc_M_N{}))> CGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
Definition: gridwise_gemm_wmma.hpp:580
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched, AEnableLds, BEnableLds >())> GridwiseGemmPipe
Definition: gridwise_gemm_wmma.hpp:148
static constexpr auto I6
Definition: gridwise_gemm_wmma.hpp:131
static __device__ void Run(const ADataType *__restrict__ p_a_grid, const BDataType *__restrict__ p_b_grid, CDataType *__restrict__ p_c_grid, void *__restrict__ p_shared, const AGridDesc &a_grid_desc, const BGridDesc &b_grid_desc, const CGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock &c_grid_desc_mblock_mperblock_nblock_nperblock, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CElementwiseOperation &c_element_op, const Block2CTileMap &block_2_ctile_map)
Definition: gridwise_gemm_wmma.hpp:585
static constexpr auto I5
Definition: gridwise_gemm_wmma.hpp:130
__host__ static constexpr __device__ auto MakeBBlockDescriptor()
Definition: gridwise_gemm_wmma.hpp:199
__host__ static constexpr __device__ auto MakeBWaveDescriptor(const BBlockDesc_ &)
Definition: gridwise_gemm_wmma.hpp:354
__host__ static constexpr __device__ bool CheckValidity(const AGridDesc &a_grid_desc, const BGridDesc &b_grid_desc, const CGridDesc_M_N &c_grid_desc_m_n, const Block2CTileMap &block_2_ctile_map)
Definition: gridwise_gemm_wmma.hpp:413
__host__ static constexpr __device__ auto MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(const CGridDesc_M_N &c_grid_desc_m_n)
Definition: gridwise_gemm_wmma.hpp:521
__host__ static constexpr __device__ auto GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat()
Definition: gridwise_gemm_wmma.hpp:399
__host__ static constexpr __device__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_wmma.hpp:513
static constexpr auto I7
Definition: gridwise_gemm_wmma.hpp:132
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_wmma.hpp:141
__host__ static constexpr __device__ auto MakeBBlockSliceCopyStep()
Definition: gridwise_gemm_wmma.hpp:268
static constexpr auto K1
Definition: gridwise_gemm_wmma.hpp:135
static constexpr auto I4
Definition: gridwise_gemm_wmma.hpp:129
static constexpr auto I1
Definition: gridwise_gemm_wmma.hpp:126
static constexpr auto MWaves
Definition: gridwise_gemm_wmma.hpp:137
__host__ static constexpr __device__ auto MakeABlockSliceCopyStep()
Definition: gridwise_gemm_wmma.hpp:248
static constexpr auto I2
Definition: gridwise_gemm_wmma.hpp:127
remove_cvref_t< decltype(MakeDefaultBlock2CTileMap(CGridDesc_M_N{}, 1, 1))> DefaultBlock2CTileMap
Definition: gridwise_gemm_wmma.hpp:582
__host__ static constexpr __device__ auto MakeAWaveDescriptor(const ABlockDesc_ &)
Definition: gridwise_gemm_wmma.hpp:290
static constexpr auto I3
Definition: gridwise_gemm_wmma.hpp:128
__host__ static constexpr __device__ auto MakeABlockDescriptor()
Definition: gridwise_gemm_wmma.hpp:151
static constexpr auto NWaves
Definition: gridwise_gemm_wmma.hpp:138
static constexpr auto I0
Definition: gridwise_gemm_wmma.hpp:125
static constexpr auto WmmaK
Definition: gridwise_gemm_wmma.hpp:139
Definition: sequence.hpp:43
Definition: tensor_space_filling_curve.hpp:20
Blockwise data transfer.
Definition: thread_group_tensor_slice_transfer_v4r1.hpp:46
Definition: thread_group_tensor_slice_transfer_v6r1.hpp:34
Definition: threadwise_tensor_slice_transfer.hpp:39
Helper structure that facilitates transfer of source (grid) data to destination threads.
Definition: threadwise_tensor_slice_transfer.hpp:234
Definition: integral_constant.hpp:20
Definition: is_known_at_compile_time.hpp:14
Definition: functional2.hpp:33
Definition: unary_element_wise_operation.hpp:334
#define CK_ENV(name)
Definition: env.hpp:129