22 template <
typename GridwiseGemm,
23 bool HasMainKBlockLoop,
25 typename Block2CTileMap,
26 typename AElementwiseOperation,
27 typename BElementwiseOperation,
28 typename CElementwiseOperation>
30 #if CK_USE_LAUNCH_BOUNDS
34 const Block2CTileMap& b2c_map,
35 const AElementwiseOperation a_element_op,
36 const BElementwiseOperation b_element_op,
37 const CElementwiseOperation c_element_op)
39 #if(defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__))
40 constexpr
index_t shared_size = GridwiseGemm::GetSharedMemoryNumberOfByte();
42 __shared__
uint8_t p_shared[shared_size];
44 GridwiseGemm::template Run<HasMainKBlockLoop, CGlobalMemoryDataOperation>(
45 karg,
static_cast<void*
>(p_shared), b2c_map, a_element_op, b_element_op, c_element_op);
63 typename AElementwiseOperation,
64 typename BElementwiseOperation,
65 typename CElementwiseOperation,
76 typename ABlockTransferThreadClusterLengths_K0_M_K1,
77 typename ABlockTransferThreadClusterArrangeOrder,
78 typename ABlockTransferSrcAccessOrder,
79 index_t ABlockTransferSrcVectorDim,
80 index_t ABlockTransferSrcScalarPerVector,
81 index_t ABlockTransferDstScalarPerVector_K1,
82 bool AThreadTransferSrcResetCoordinateAfterRun,
84 typename BBlockTransferThreadClusterLengths_K0_N_K1,
85 typename BBlockTransferThreadClusterArrangeOrder,
86 typename BBlockTransferSrcAccessOrder,
87 index_t BBlockTransferSrcVectorDim,
88 index_t BBlockTransferSrcScalarPerVector,
89 index_t BBlockTransferDstScalarPerVector_K1,
90 bool BThreadTransferSrcResetCoordinateAfterRun,
92 index_t CShuffleMRepeatPerShuffle,
93 index_t CShuffleNRepeatPerShuffle,
94 index_t CBlockTransferScalarPerVector_NWaveNPerXDL,
95 typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
98 typename ComputeTypeA = FloatC,
99 typename ComputeTypeB = ComputeTypeA,
100 typename LDSTypeA = ComputeTypeA,
101 typename LDSTypeB = ComputeTypeB>
115 static constexpr
auto M01 = 1;
116 static constexpr
auto N01 = 1;
120 MPerBlock, NPerBlock,
K1* K0PerBlock};
125 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
145 const FloatB* p_b_grid_,
177 std::cout <<
"arg {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
181 <<
"KB:" <<
k_batch <<
"}" << std::endl;
206 auto K_t = K_Batch * K0PerBlock *
K1;
207 return (K + K_t - 1) / K_t * K0PerBlock;
213 return K_Batch * K0Padded *
K1;
224 const auto a_grid_desc_m_k = [&]() {
300 const auto b_grid_desc_k_n = [&]() {
370 const auto c_grid_desc_m_n = [&]() {
381 return gemm_padder.PadCDescriptor_M_N(c_grid_desc_m_n);
386 constexpr
auto max_lds_align =
K1;
389 constexpr
auto a_k0_m_k1_block_desc = [&]() {
390 if constexpr(ABlockLdsExtraM)
404 constexpr
auto b_k0_n_k1_block_desc = [&]() {
405 if constexpr(BBlockLdsExtraN)
419 constexpr
auto a_block_space_size =
422 constexpr
auto b_block_space_size =
425 constexpr
auto c_block_size =
428 return math::max(a_block_space_size *
sizeof(LDSTypeA) +
429 b_block_space_size *
sizeof(LDSTypeB),
430 c_block_size *
sizeof(FloatC));
440 if(!(karg.
M % MPerBlock == 0))
444 std::cout <<
"Arg M value is not a multiple of MPerBlock! M: " << karg.
M <<
" "
445 << __FILE__ <<
":" << __LINE__ <<
", in function: " << __func__
457 if(!(karg.
N % NPerBlock == 0))
461 std::cout <<
"Arg N value is not a multiple of NPerBlock! N: " << karg.
N <<
" "
462 << __FILE__ <<
":" << __LINE__ <<
", in function: " << __func__
475 auto K_t = karg.
k_batch * K0PerBlock *
K1;
476 if(!(karg.
K % K_t == 0))
480 std::cout <<
"Arg K value is not a multiple of K_Batch * K0PerBlock * K1! K: "
481 << karg.
K <<
" " << __FILE__ <<
":" << __LINE__
482 <<
", in function: " << __func__ << std::endl;
490 if(karg.
K % ABlockTransferSrcScalarPerVector != 0)
494 std::cout <<
"Arg K (" << karg.
K
495 <<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
496 << ABlockTransferSrcScalarPerVector <<
" )! " << __FILE__ <<
":"
497 << __LINE__ <<
", in function: " << __func__ << std::endl;
504 if(karg.
M % ABlockTransferSrcScalarPerVector != 0)
508 std::cout <<
"Arg M (" << karg.
M
509 <<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
510 << ABlockTransferSrcScalarPerVector <<
" )! " << __FILE__ <<
":"
511 << __LINE__ <<
", in function: " << __func__ << std::endl;
519 if(karg.
N % BBlockTransferSrcScalarPerVector != 0)
523 std::cout <<
"Arg N (" << karg.
N
524 <<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
525 << BBlockTransferSrcScalarPerVector <<
" )! " << __FILE__ <<
":"
526 << __LINE__ <<
", in function: " << __func__ << std::endl;
533 if(karg.
K % BBlockTransferSrcScalarPerVector != 0)
537 std::cout <<
"Arg K (" << karg.
K
538 <<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
539 << BBlockTransferSrcScalarPerVector <<
" )! " << __FILE__ <<
":"
540 << __LINE__ <<
", in function: " << __func__ << std::endl;
548 if(karg.
N % CBlockTransferScalarPerVector_NWaveNPerXDL != 0)
552 std::cout <<
"Arg N (" << karg.
N
553 <<
") value is not a multiple of "
554 "CBlockTransferScalarPerVector_NWaveNPerXDL ("
555 << CBlockTransferScalarPerVector_NWaveNPerXDL <<
" )! " << __FILE__
556 <<
":" << __LINE__ <<
", in function: " << __func__ << std::endl;
563 if(karg.
M % CBlockTransferScalarPerVector_NWaveNPerXDL != 0)
567 std::cout <<
"Arg M (" << karg.
M
568 <<
") value is not a multiple of "
569 "CBlockTransferScalarPerVector_NWaveNPerXDL ("
570 << CBlockTransferScalarPerVector_NWaveNPerXDL <<
" )! " << __FILE__
571 <<
":" << __LINE__ <<
", in function: " << __func__ << std::endl;
577 const auto num_k_loop = karg.
K0Padded / K0PerBlock;
578 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
582 std::cout <<
"The number of k loops (" << num_k_loop
583 <<
") value is not supported by GridwiseGemm Pipeline."
584 <<
" K0Padded: " << karg.
K0Padded <<
", K0PerBlock: " << K0PerBlock <<
" "
585 << __FILE__ <<
":" << __LINE__ <<
", in function: " << __func__
598 const index_t KPad = KBatch * K0Padded *
K1;
604 const index_t num_loop = K0Padded / K0PerBlock;
605 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
608 template <
typename CGr
idDesc>
609 __host__ __device__
static constexpr
auto
612 const auto M = c_m_n_grid_desc.GetLength(
I0);
613 const auto N = c_m_n_grid_desc.GetLength(
I1);
615 const auto MBlock = M / MPerBlock;
616 const auto NBlock = N / NPerBlock;
627 template <
typename CGr
idDesc>
632 c_m_n_grid_desc, 8, KBatch);
635 __host__ __device__
static constexpr
auto
638 constexpr
index_t MWave = MPerBlock / (MRepeat * MPerXDL);
639 constexpr
index_t NWave = NPerBlock / (NRepeat * NPerXDL);
657 template <
bool HasMainKBlockLoop,
659 typename Block2CTileMap>
661 void* __restrict__ p_shared_block,
662 const Block2CTileMap& block_2_ctile_map,
663 const AElementwiseOperation a_element_op = AElementwiseOperation{},
664 const BElementwiseOperation b_element_op = BElementwiseOperation{},
665 const CElementwiseOperation c_element_op = CElementwiseOperation{})
667 const FloatA* p_a_grid = karg.
p_a_grid;
668 const FloatB* p_b_grid = karg.
p_b_grid;
676 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
679 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
680 p_a_grid, a_b_k0_m_k1_grid_desc.GetElementSpaceSize());
681 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
682 p_b_grid, b_b_k0_n_k1_grid_desc.GetElementSpaceSize());
683 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
684 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
687 const auto block_work_idx =
690 if(!block_2_ctile_map.ValidCTileIndex(
692 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
693 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
698 const index_t block_m_id = __builtin_amdgcn_readfirstlane(block_work_idx[
I1]);
699 const index_t block_n_id = __builtin_amdgcn_readfirstlane(block_work_idx[
I2]);
700 const index_t k_batch_id = __builtin_amdgcn_readfirstlane(block_work_idx[
I0]);
703 const index_t m_block_data_idx_on_grid =
704 __builtin_amdgcn_readfirstlane(block_m_id * MPerBlock);
706 const index_t n_block_data_idx_on_grid =
707 __builtin_amdgcn_readfirstlane(block_n_id * NPerBlock);
710 constexpr
auto max_lds_align =
K1;
713 constexpr
auto a_k0_m_k1_block_desc = [&]() {
714 if constexpr(ABlockLdsExtraM)
717 make_tuple(Number<K0PerBlock>{}, Number<MPerBlock>{},
K1),
723 make_tuple(Number<K0PerBlock>{}, Number<MPerBlock>{},
K1), max_lds_align);
727 constexpr
auto a_b_k0_m_k1_block_desc = [&]() {
728 if constexpr(ABlockLdsExtraM)
731 make_tuple(Number<1>{}, Number<K0PerBlock>{}, Number<MPerBlock>{},
K1),
732 make_tuple(Number<K0PerBlock>{} * Number<MPerBlock + 1>{} *
K1,
733 Number<MPerBlock + 1>{} *
K1,
740 make_tuple(Number<1>{}, Number<K0PerBlock>{}, Number<MPerBlock>{},
K1),
745 constexpr
auto b_k0_n_k1_block_desc = [&]() {
746 if constexpr(BBlockLdsExtraN)
749 make_tuple(Number<K0PerBlock>{}, Number<NPerBlock>{},
K1),
755 make_tuple(Number<K0PerBlock>{}, Number<NPerBlock>{},
K1), max_lds_align);
759 constexpr
auto b_b_k0_n_k1_block_desc = [&]() {
760 if constexpr(BBlockLdsExtraN)
763 make_tuple(Number<1>{}, Number<K0PerBlock>{}, Number<NPerBlock>{},
K1),
764 make_tuple(Number<K0PerBlock>{} * Number<NPerBlock + 1>{} *
K1,
765 Number<NPerBlock + 1>{} *
K1,
772 make_tuple(Number<1>{}, Number<K0PerBlock>{}, Number<NPerBlock>{},
K1),
777 auto a_blockwise_copy =
779 AElementwiseOperation,
782 Sequence<1, K0PerBlock, MPerBlock, K1>,
783 ABlockTransferThreadClusterLengths_K0_M_K1,
784 ABlockTransferThreadClusterArrangeOrder,
787 decltype(a_b_k0_m_k1_grid_desc),
788 decltype(a_b_k0_m_k1_block_desc),
789 ABlockTransferSrcAccessOrder,
790 Sequence<0, 2, 1, 3>,
791 ABlockTransferSrcVectorDim,
793 ABlockTransferSrcScalarPerVector,
794 ABlockTransferDstScalarPerVector_K1,
797 AThreadTransferSrcResetCoordinateAfterRun,
799 a_b_k0_m_k1_grid_desc,
802 a_b_k0_m_k1_block_desc,
807 auto b_blockwise_copy =
809 BElementwiseOperation,
812 Sequence<1, K0PerBlock, NPerBlock, K1>,
813 BBlockTransferThreadClusterLengths_K0_N_K1,
814 BBlockTransferThreadClusterArrangeOrder,
817 decltype(b_b_k0_n_k1_grid_desc),
818 decltype(b_b_k0_n_k1_block_desc),
819 BBlockTransferSrcAccessOrder,
820 Sequence<0, 2, 1, 3>,
821 BBlockTransferSrcVectorDim,
823 BBlockTransferSrcScalarPerVector,
824 BBlockTransferDstScalarPerVector_K1,
827 BThreadTransferSrcResetCoordinateAfterRun,
829 b_b_k0_n_k1_grid_desc,
832 b_b_k0_n_k1_block_desc,
849 decltype(a_k0_m_k1_block_desc),
850 decltype(b_k0_n_k1_block_desc),
860 auto c_thread_buf = blockwise_gemm.GetCThreadBuffer();
863 constexpr
auto a_block_space_size =
866 auto p_a_block =
reinterpret_cast<LDSTypeA*
>(p_shared_block);
867 auto p_b_block =
reinterpret_cast<LDSTypeB*
>(p_a_block + a_block_space_size);
869 constexpr
auto a_block_slice_copy_step =
make_multi_index(0, K0PerBlock, 0, 0);
870 constexpr
auto b_block_slice_copy_step =
make_multi_index(0, K0PerBlock, 0, 0);
872 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
873 p_a_block, a_k0_m_k1_block_desc.GetElementSpaceSize());
874 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
875 p_b_block, b_k0_n_k1_block_desc.GetElementSpaceSize());
878 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
879 (a_b_k0_m_k1_grid_desc.GetLength(
I1) * a_b_k0_m_k1_grid_desc.GetLength(
I3)) /
884 gridwise_gemm_pipeline.template Run<HasMainKBlockLoop>(a_b_k0_m_k1_grid_desc,
885 a_b_k0_m_k1_block_desc,
889 a_block_slice_copy_step,
890 b_b_k0_n_k1_grid_desc,
891 b_b_k0_n_k1_block_desc,
895 b_block_slice_copy_step,
898 num_k_block_main_loop);
902 constexpr
index_t MWave = MPerBlock / (MRepeat * MPerXDL);
903 constexpr
index_t NWave = NPerBlock / (NRepeat * NPerXDL);
905 constexpr
auto c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc =
906 blockwise_gemm.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
908 constexpr
auto c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc =
909 blockwise_gemm.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
911 constexpr
auto M0 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I0);
912 constexpr
auto N0 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I1);
913 constexpr
auto M1 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I2);
914 constexpr
auto N1 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I3);
915 constexpr
auto M2 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I4);
916 constexpr
auto M3 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I5);
917 constexpr
auto M4 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I6);
918 constexpr
auto N2 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(
I7);
920 constexpr
auto c_block_desc_mblock_mperblock_nblock_nperblock =
923 auto c_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
924 static_cast<FloatC*
>(p_shared_block),
925 c_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
928 c_block_desc_mblock_mperblock_nblock_nperblock,
940 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
942 Sequence<>{}, Sequence<0, 2, 4, 5, 6>{}, Sequence<>{}, Sequence<1, 3, 7>{}));
946 const auto c_thread_mtx_on_block =
947 blockwise_gemm.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
949 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
950 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
952 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
958 const auto m_thread_data_on_block_idx =
959 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
962 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
968 const auto n_thread_data_on_block_idx =
969 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
973 auto c_thread_copy_vgpr_to_lds =
974 ThreadwiseTensorSliceTransfer_v1r3<FloatAcc,
976 decltype(c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc),
977 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
979 Sequence<CShuffleMRepeatPerShuffle,
980 CShuffleNRepeatPerShuffle,
987 Sequence<0, 1, 2, 3, 4, 5, 6, 7>,
993 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
996 m_thread_data_on_block_idx[
I1],
997 n_thread_data_on_block_idx[
I1],
998 m_thread_data_on_block_idx[
I2],
999 m_thread_data_on_block_idx[
I3],
1000 m_thread_data_on_block_idx[
I4],
1001 n_thread_data_on_block_idx[
I2]),
1005 auto c_block_copy_lds_to_global = ThreadGroupTensorSliceTransfer_v6r1<
1007 CElementwiseOperation,
1008 CGlobalMemoryDataOperation,
1010 CShuffleMRepeatPerShuffle * MWave * MPerXDL,
1012 CShuffleNRepeatPerShuffle * NWave * NPerXDL>,
1013 CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1014 Sequence<0, 1, 2, 3>,
1017 decltype(c_block_desc_mblock_mperblock_nblock_nperblock),
1018 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1019 Sequence<0, 1, 2, 3>,
1021 CBlockTransferScalarPerVector_NWaveNPerXDL,
1024 {c_block_desc_mblock_mperblock_nblock_nperblock,
1026 c_grid_desc_mblock_mperblock_nblock_nperblock,
1030 constexpr
auto mxdlperwave_forward_step =
1032 constexpr
auto nxdlperwave_forward_step =
1034 constexpr
auto nxdlperwave_backward_step =
1037 static_for<0, MRepeat, CShuffleMRepeatPerShuffle>{}([&](
auto mxdlperwave_iter) {
1038 constexpr
auto mxdlperwave = mxdlperwave_iter;
1040 static_for<0, NRepeat, CShuffleNRepeatPerShuffle>{}([&](
auto nxdlperwave_iter) {
1041 constexpr
bool nxdlperwave_forward_sweep =
1042 (mxdlperwave % (2 * CShuffleMRepeatPerShuffle) == 0);
1044 constexpr
index_t nxdlperwave_value =
1045 nxdlperwave_forward_sweep
1047 : (NRepeat - nxdlperwave_iter - CShuffleNRepeatPerShuffle);
1049 constexpr
auto nxdlperwave = Number<nxdlperwave_value>{};
1055 c_thread_copy_vgpr_to_lds.Run(
1056 c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc,
1059 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1066 c_block_copy_lds_to_global.Run(c_block_desc_mblock_mperblock_nblock_nperblock,
1068 c_grid_desc_mblock_mperblock_nblock_nperblock,
1072 if constexpr(nxdlperwave_forward_sweep &&
1073 (nxdlperwave < NRepeat - CShuffleNRepeatPerShuffle))
1075 c_block_copy_lds_to_global.MoveDstSliceWindow(
1076 c_grid_desc_mblock_mperblock_nblock_nperblock,
1077 nxdlperwave_forward_step);
1079 else if constexpr((!nxdlperwave_forward_sweep) && (nxdlperwave > 0))
1081 c_block_copy_lds_to_global.MoveDstSliceWindow(
1082 c_grid_desc_mblock_mperblock_nblock_nperblock,
1083 nxdlperwave_backward_step);
1088 if constexpr(mxdlperwave < MRepeat - CShuffleMRepeatPerShuffle)
1090 c_block_copy_lds_to_global.MoveDstSliceWindow(
1091 c_grid_desc_mblock_mperblock_nblock_nperblock, mxdlperwave_forward_step);
1099 auto str = std::stringstream();
1102 str <<
"GemmXdlSplitKCShuffle_"
1104 << std::string(ALayout::name)[0]
1105 << std::string(BLayout::name)[0]
1106 << std::string(CLayout::name)[0]
1108 <<
"B" << BlockSize <<
"_"
1109 <<
"Vec" << ABlockTransferSrcScalarPerVector <<
"x"
1110 << BBlockTransferSrcScalarPerVector <<
"x"
1111 << CBlockTransferScalarPerVector_NWaveNPerXDL <<
"_"
1114 << K0PerBlock <<
"x"
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:30
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:29
__host__ constexpr __device__ auto integer_least_multiple(X x, Y y)
Definition: math.hpp:78
__host__ constexpr __device__ auto integer_divide_ceil(X x, Y y)
Definition: math.hpp:72
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
std::string getGemmSpecializationString(const GemmSpecialization &s)
Definition: gemm_specialization.hpp:32
GemmSpecialization
Definition: gemm_specialization.hpp:11
__host__ constexpr __device__ auto make_multi_index(Xs &&... xs)
Definition: array_multi_index.hpp:15
constexpr auto BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_Selector()
Definition: blockwise_gemm_xdlops.hpp:605
__host__ constexpr __device__ auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition: tensor_descriptor_helper.hpp:49
__global__ void kernel_gemm_xdlops_v2r4r2_simplified(typename GridwiseGemm::Argument karg, const Block2CTileMap &b2c_map, const AElementwiseOperation a_element_op, const BElementwiseOperation b_element_op, const CElementwiseOperation c_element_op)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:33
InMemoryDataOperationEnum
Definition: ck.hpp:276
__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_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:98
constexpr detail::ignore_t ignore
Definition: ignore.hpp:20
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:58
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:298
__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_descriptor.hpp:319
__host__ constexpr __device__ auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:37
__device__ void block_sync_lds()
Definition: synchronization.hpp:10
PipelineVersion
Definition: gridwise_gemm_pipeline_selector.hpp:18
constexpr LoopScheduler make_default_loop_scheduler()
Definition: loop_scheduler.hpp:20
unsigned char uint8_t
Definition: stdint.h:124
Simple tile mapping which creates 3D grid of block of threads.
Definition: block_to_ctile_map.hpp:976
Definition: block_to_ctile_map.hpp:540
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:128
index_t StrideC
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:137
index_t K0Padded
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:141
index_t MPadded
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:138
index_t k_batch
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:142
index_t N
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:133
const FloatA * p_a_grid
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:129
index_t K
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:134
index_t StrideA
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:135
index_t M
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:132
const FloatB * p_b_grid
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:130
index_t StrideB
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:136
FloatC * p_c_grid
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:131
index_t NPadded
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:139
index_t KPadded
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:140
Argument(const FloatA *p_a_grid_, const FloatB *p_b_grid_, FloatC *p_c_grid_, index_t M_, index_t N_, index_t K_, index_t StrideA_, index_t StrideB_, index_t StrideC_, index_t MPadded_, index_t NPadded_, index_t KPadded_, index_t K0Padded_, index_t k_batch_)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:144
void Print() const
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:175
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:103
static constexpr auto I2
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:106
static constexpr auto gemm_padder
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:118
__host__ static __device__ auto CalculateGridSize(const Argument &karg)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:185
__host__ static __device__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:193
static constexpr auto I5
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:109
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:125
__host__ static __device__ auto GetKPad(index_t K, index_t KBatch)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:594
remove_cvref_t< decltype(MakeDefaultBlock2CTileMap())> DefaultBlock2CTileMap
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:655
static constexpr auto K1
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:114
static constexpr auto I1
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:105
static __device__ void Run(const Argument &karg, void *__restrict__ p_shared_block, const Block2CTileMap &block_2_ctile_map, const AElementwiseOperation a_element_op=AElementwiseOperation{}, const BElementwiseOperation b_element_op=BElementwiseOperation{}, const CElementwiseOperation c_element_op=CElementwiseOperation{})
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:660
__host__ static constexpr __device__ auto MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(const CGridDesc &c_m_n_grid_desc)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:610
__host__ static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:384
__host__ static constexpr __device__ bool CheckValidity(const Argument &karg)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:433
static constexpr auto I3
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:107
static constexpr auto I7
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:111
__host__ static constexpr __device__ auto GetCBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:636
static constexpr auto I6
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:110
__host__ static __device__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:198
static constexpr auto N01
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:116
__host__ static __device__ auto MakeAGridDescriptor_KBatch_K0_M_K1(index_t M, index_t MPad, index_t K, index_t StrideA, index_t KBatch, index_t K0Padded, index_t KPad)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:216
__host__ static __device__ auto CalculateK0Padded(index_t K, index_t K_Batch=1)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:203
__host__ static constexpr __device__ bool CalculateHasMainK0BlockLoop(index_t K0Padded)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:602
__host__ static constexpr __device__ auto MakeDefaultBlock2CTileMap()
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:649
__host__ static __device__ auto MakeBGridDescriptor_KBatch_K0_N_K1(index_t K, index_t NPad, index_t N, index_t StrideB, index_t KBatch, index_t K0Padded, index_t KPad)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:292
__host__ static constexpr __device__ auto MakeCBlockClusterAdaptor(const CGridDesc &c_m_n_grid_desc, index_t, index_t, index_t KBatch)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:628
__host__ static __device__ auto MakeCGridDescriptor_M_N(index_t M, index_t N, index_t StrideC)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:368
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:122
static constexpr auto M01
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:115
remove_cvref_t< decltype(MakeCGridDescriptor_M_N(1, 1, 1))> CGridDesc_M_N
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:654
__host__ static __device__ auto CalculateKPadded(index_t K, index_t K_Batch=1)
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:210
static std::string GetTypeString()
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:1097
static constexpr auto I4
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:108
static constexpr auto I0
Definition: gridwise_gemm_xdlops_v2r4r2.hpp:104
Definition: sequence.hpp:43
Definition: integral_constant.hpp:20
Definition: device_base.hpp:51
Definition: matrix_padder.hpp:134
Definition: unary_element_wise_operation.hpp:334
#define CK_ENV(name)
Definition: env.hpp:129