20 template <
typename Gr
idwiseGemm,
bool HasMainKBlockLoop, index_t TailNum = 3>
22 #if CK_USE_LAUNCH_BOUNDS
28 #if(defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__))
31 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
32 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
34 GridwiseGemm::template Run<HasMainKBlockLoop, TailNum>(
35 karg.p_a_grid, karg.p_b_grid, karg.p_c_grid, p_shared_0, p_shared_1, karg);
41 template <
typename GridwiseGemm,
45 bool HasMainKBlockLoop>
47 #if CK_USE_LAUNCH_BOUNDS
51 const FloatB* p_b_grid,
53 typename GridwiseGemm::Problem problem)
55 #if(defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__))
56 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
57 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
59 GridwiseGemm::template Run<HasMainKBlockLoop>(
60 p_a_grid, p_b_grid, p_c_grid, p_shared_0, p_shared_1, problem);
69 template <
typename ALayout,
74 typename FloatGemmAcc,
75 typename FloatCShuffle,
77 typename AElementwiseOperation,
78 typename BElementwiseOperation,
79 typename CElementwiseOperation,
93 typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
94 typename ABlockTransferThreadClusterArrangeOrder,
95 typename ABlockTransferSrcAccessOrder,
96 index_t ABlockTransferSrcVectorDim,
97 index_t ABlockTransferSrcScalarPerVector,
98 index_t ABlockTransferDstScalarPerVector_AK1,
99 bool AThreadTransferSrcResetCoordinateAfterRun,
101 typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
102 typename BBlockTransferThreadClusterArrangeOrder,
103 typename BBlockTransferSrcAccessOrder,
104 index_t BBlockTransferSrcVectorDim,
105 index_t BBlockTransferSrcScalarPerVector,
106 index_t BBlockTransferDstScalarPerVector_BK1,
107 bool BThreadTransferSrcResetCoordinateAfterRun,
109 index_t CShuffleMXdlPerWavePerShuffle,
110 index_t CShuffleNXdlPerWavePerShuffle,
111 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
112 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
115 typename ComputeTypeA = FloatC,
116 typename ComputeTypeB = ComputeTypeA>
160 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
161 GemmSpec == GemmSpecialization::MNKPadding ||
162 GemmSpec == GemmSpecialization::KPadding ||
163 GemmSpec == GemmSpecialization::NKPadding)
177 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
178 GemmSpec == GemmSpecialization::MNKPadding ||
179 GemmSpec == GemmSpecialization::KPadding ||
180 GemmSpec == GemmSpecialization::MKPadding)
200 template <index_t MNXdlPerWave, index_t MNWaves, index_t MNPerXdl,
typename TileDesc_K0_MN_K1>
218 const auto a_grid_desc_mraw_kraw = [&]() {
219 if constexpr(is_same_v<tensor_layout::gemm::RowMajor, ALayout>)
223 else if constexpr(is_same_v<tensor_layout::gemm::ColumnMajor, ALayout>)
231 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
232 GemmSpec == GemmSpecialization::MNKPadding)
235 const auto a_grid_desc_m_k =
249 return a_grid_desc_ak0_m_ak1;
251 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
252 GemmSpec == GemmSpecialization::MNPadding)
256 a_grid_desc_mraw_kraw,
262 return a_grid_desc_ak0_m_ak1;
264 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
265 GemmSpec == GemmSpecialization::NKPadding)
269 a_grid_desc_mraw_kraw,
281 return a_grid_desc_ak0_m_ak1;
287 a_grid_desc_mraw_kraw,
293 return a_grid_desc_ak0_m_ak1;
300 const auto b_grid_desc_nraw_kraw = [&]() {
313 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
314 GemmSpec == GemmSpecialization::MNKPadding)
317 const auto b_grid_desc_n_k =
331 return b_grid_desc_bk0_n_bk1;
333 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
334 GemmSpec == GemmSpecialization::MNPadding)
338 b_grid_desc_nraw_kraw,
344 return b_grid_desc_bk0_n_bk1;
346 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
347 GemmSpec == GemmSpecialization::MKPadding)
351 b_grid_desc_nraw_kraw,
363 return b_grid_desc_bk0_n_bk1;
369 b_grid_desc_nraw_kraw,
375 return b_grid_desc_bk0_n_bk1;
379 template <
typename ABlockDesc_AK0_M_AK1>
380 __host__ __device__
static constexpr
auto
383 constexpr
index_t MWaves = MPerBlock / (MXdlPerWave * MPerXdl);
385 return MakeGemmMmaTileDescriptor<MXdlPerWave, MWaves, MPerXdl>(ABlockDesc_AK0_M_AK1{});
388 template <
typename BBlockDesc_BK0_N_BK1>
389 __host__ __device__
static constexpr
auto
392 constexpr
index_t NWaves = NPerBlock / (NXdlPerWave * NPerXdl);
394 return MakeGemmMmaTileDescriptor<NXdlPerWave, NWaves, NPerXdl>(BBlockDesc_BK0_N_BK1{});
397 __host__ __device__
static auto
400 const auto c_grid_desc_mraw_nraw = [&]() {
413 if constexpr(GemmSpec == GemmSpecialization::MNPadding ||
414 GemmSpec == GemmSpecialization::MNKPadding)
423 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
424 GemmSpec == GemmSpecialization::MKPadding)
428 c_grid_desc_mraw_nraw,
433 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
434 GemmSpec == GemmSpecialization::NKPadding)
438 c_grid_desc_mraw_nraw,
446 return c_grid_desc_mraw_nraw;
476 std::cout <<
"problem {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
479 <<
"KP:" <<
KPadded <<
", " <<
"AK0:" <<
AK0 <<
", " <<
"BK0:" <<
BK0 <<
", "
480 <<
"MBlock: " <<
MBlock <<
", " <<
"NBlock: " <<
NBlock <<
"}" << std::endl;
502 const FloatB* p_b_grid_,
510 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
524 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
544 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
545 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
547 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
554 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock;
567 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
570 b_block_desc_bk0_n_bk1.GetElementSpaceSize(), max_lds_align);
573 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
576 constexpr
auto c_block_size =
577 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize();
579 return math::max((a_block_space_size_aligned *
sizeof(ComputeTypeA) +
580 b_block_space_size_aligned *
sizeof(ComputeTypeB)),
581 c_block_size *
sizeof(FloatCShuffle));
587 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
588 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
589 "Invalid tuning param!");
596 if(!(problem.
M % MPerBlock == 0))
607 if(!(problem.
N % NPerBlock == 0))
626 if(!(problem.
K % AK1Value == 0) || !(problem.
K % BK1Value == 0))
634 if(problem.
K % ABlockTransferSrcScalarPerVector != 0)
641 if(problem.
M % ABlockTransferSrcScalarPerVector != 0)
649 if(problem.
N % BBlockTransferSrcScalarPerVector != 0)
656 if(problem.
K % BBlockTransferSrcScalarPerVector != 0)
664 if(problem.
N % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
671 if(problem.
M % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
678 const auto num_k_loop = (
CalculateAK0(problem.
K) * AK1Value) / KPerBlock;
691 const index_t num_loop = K / KPerBlock;
698 const index_t num_loop = K / KPerBlock;
700 if(num_loop % 2 == 1)
706 template <
typename CGr
idDesc>
717 return c_grid_desc_mblock_mperblock_nblock_nperblock;
724 template <
bool HasMainKBlockLoop, index_t TailNum = 3>
725 __device__
static void Run(
const FloatA* p_a_grid,
726 const FloatB* p_b_grid,
739 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
743 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
744 p_a_grid, a_grid_desc_ak0_m_ak1.GetElementSpaceSize());
745 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
746 p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize());
747 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
748 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
750 const AElementwiseOperation a_element_op{};
751 const BElementwiseOperation b_element_op{};
752 const CElementwiseOperation c_element_op{};
757 const auto block_work_idx =
760 if(!block_2_ctile_map.ValidCTileIndex(
762 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
763 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
768 if(threadIdx.x == 0){
769 printf(
"Hardware assigned No. %03d workgroup of logical C tile (%02d, %02d) on %d th XCC Die, %d th SE, %d th CU\n",
779 const index_t m_block_data_idx_on_grid =
780 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
782 const index_t n_block_data_idx_on_grid =
783 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
795 auto a_blockwise_copy =
797 AElementwiseOperation,
801 ABlockTransferThreadClusterLengths_AK0_M_AK1,
802 ABlockTransferThreadClusterArrangeOrder,
805 decltype(a_grid_desc_ak0_m_ak1),
806 decltype(a_block_desc_ak0_m_ak1),
807 ABlockTransferSrcAccessOrder,
809 ABlockTransferSrcVectorDim,
811 ABlockTransferSrcScalarPerVector,
812 ABlockTransferDstScalarPerVector_AK1,
815 AThreadTransferSrcResetCoordinateAfterRun,
817 a_grid_desc_ak0_m_ak1,
820 a_block_desc_ak0_m_ak1,
825 auto b_blockwise_copy =
827 BElementwiseOperation,
831 BBlockTransferThreadClusterLengths_BK0_N_BK1,
832 BBlockTransferThreadClusterArrangeOrder,
835 decltype(b_grid_desc_bk0_n_bk1),
836 decltype(b_block_desc_bk0_n_bk1),
837 BBlockTransferSrcAccessOrder,
839 BBlockTransferSrcVectorDim,
841 BBlockTransferSrcScalarPerVector,
842 BBlockTransferDstScalarPerVector_BK1,
845 BThreadTransferSrcResetCoordinateAfterRun,
847 b_grid_desc_bk0_n_bk1,
850 b_block_desc_bk0_n_bk1,
862 constexpr
bool is_single_rate_mfma =
870 constexpr
auto is_scale_mfma =
false;
877 is_scale_mfma>::selected_mfma.k_per_blk);
895 decltype(a_block_desc_ak0_m_ak1),
896 decltype(b_block_desc_bk0_n_bk1),
908 auto c_thread_buf = blockwise_gemm_pipeline.GetCThreadBuffer();
912 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
914 auto a_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
915 static_cast<ComputeTypeA*
>(p_shared_0), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
917 auto b_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
918 static_cast<ComputeTypeB*
>(p_shared_0) + a_block_space_size_aligned,
919 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
921 auto a_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
922 static_cast<ComputeTypeA*
>(p_shared_1), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
924 auto b_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
925 static_cast<ComputeTypeB*
>(p_shared_1) + a_block_space_size_aligned,
926 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
928 auto a_block_bufs =
make_tuple(a_block_buf_ping, a_block_buf_pong);
929 auto b_block_bufs =
make_tuple(b_block_buf_ping, b_block_buf_pong);
935 static_assert(std::is_default_constructible_v<GridwiseGemmPipe>);
938 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
939 (a_grid_desc_ak0_m_ak1.GetLength(
I0) * a_grid_desc_ak0_m_ak1.GetLength(
I2)) /
942 blockwise_gemm_pipeline.template Run<HasMainKBlockLoop, TailNum>(a_grid_desc_ak0_m_ak1,
943 a_block_desc_ak0_m_ak1,
947 a_block_slice_copy_step,
948 b_grid_desc_bk0_n_bk1,
949 b_block_desc_bk0_n_bk1,
953 b_block_slice_copy_step,
955 num_k_block_main_loop);
959 static_assert(MXdlPerWave % CShuffleMXdlPerWavePerShuffle == 0 &&
960 NXdlPerWave % CShuffleNXdlPerWavePerShuffle == 0,
963 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
964 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
967 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
968 blockwise_gemm_pipeline.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
972 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp =
973 blockwise_gemm_pipeline.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
975 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I0);
976 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I1);
977 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I2);
978 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I3);
979 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I4);
980 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I5);
981 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I6);
982 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I7);
984 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
987 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
988 static_cast<FloatCShuffle*
>(p_shared_0),
989 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
992 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1012 const auto c_thread_mtx_on_block =
1013 blockwise_gemm_pipeline.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
1015 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
1016 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
1018 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
1024 const auto m_thread_data_on_block_idx =
1025 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
1028 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
1034 const auto n_thread_data_on_block_idx =
1035 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
1039 auto c_thread_copy_vgpr_to_lds =
1042 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1043 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1045 Sequence<CShuffleMXdlPerWavePerShuffle,
1046 CShuffleNXdlPerWavePerShuffle,
1059 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1062 m_thread_data_on_block_idx[
I1],
1063 n_thread_data_on_block_idx[
I1],
1064 m_thread_data_on_block_idx[
I2],
1065 m_thread_data_on_block_idx[
I3],
1066 m_thread_data_on_block_idx[
I4],
1067 n_thread_data_on_block_idx[
I2]),
1073 CElementwiseOperation,
1074 CGlobalMemoryDataOperation,
1076 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1078 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>,
1079 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1083 decltype(c_shuffle_block_desc_mblock_mperblock_nblock_nperblock),
1084 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1087 CShuffleBlockTransferScalarPerVector_NPerBlock,
1090 {c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1092 c_grid_desc_mblock_mperblock_nblock_nperblock,
1097 constexpr
auto sfc_c_vgpr =
1100 Sequence<CShuffleMXdlPerWavePerShuffle,
1101 CShuffleNXdlPerWavePerShuffle,
1110 constexpr
auto sfc_c_global =
1114 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1116 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{};
1118 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
1120 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1127 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1128 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1130 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1131 c_shuffle_block_buf);
1137 c_shuffle_block_copy_lds_to_global.Run(
1138 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1139 c_shuffle_block_buf,
1140 c_grid_desc_mblock_mperblock_nblock_nperblock,
1143 if constexpr(access_id < num_access - 1)
1145 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1148 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1149 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:29
Y __host__ constexpr __device__ auto lcm(X x, Y y)
Definition: math.hpp:198
__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__ auto integer_divide_floor(X x, Y y)
Definition: math.hpp:66
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
GemmSpecialization
Definition: gemm_specialization.hpp:11
__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
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_merge_transform_v3_division_mod(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:84
__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
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__global__ void kernel_gemm_xdl_cshuffle_v2(typename GridwiseGemm::Argument karg)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:26
__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
Definition: block_to_ctile_map.hpp:270
__host__ static constexpr __device__ index_t CalculateGridSize(index_t M, index_t N)
Definition: block_to_ctile_map.hpp:282
Definition: blockwise_gemm_pipeline_xdlops.hpp:103
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:500
const FloatB * p_b_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:518
__host__ 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_)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:501
FloatC * p_c_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:519
const FloatA * p_a_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:517
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:451
index_t StrideB
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:487
index_t StrideC
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:488
index_t KPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:491
index_t N
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:484
index_t M
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:483
index_t MBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:494
index_t NBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:495
__host__ void Print() const
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:474
index_t BK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:493
index_t MPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:489
index_t K
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:485
index_t StrideA
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:486
index_t AK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:492
__host__ Problem(index_t M_, index_t N_, index_t K_, index_t StrideA_, index_t StrideB_, index_t StrideC_)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:452
index_t NPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:490
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:118
static constexpr auto I6
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:125
static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:557
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:141
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:136
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:689
static constexpr __host__ index_t CalculateKBlockLoopTailNum(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:696
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:524
static constexpr __device__ auto MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(const CGridDesc &c_grid_desc_m_n, index_t MBlock, index_t NBlock)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:707
static __device__ void Run(const FloatA *p_a_grid, const FloatB *p_b_grid, FloatC *p_c_grid, void *p_shared_0, void *p_shared_1, const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:725
static constexpr auto AK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:131
static constexpr auto I1
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:120
static constexpr auto I2
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:121
static constexpr auto AK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:129
static constexpr auto I0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:119
static __host__ auto CalculateMBlock(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:190
static constexpr auto BK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:130
static __host__ auto CalculateAK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:156
static __host__ auto CalculateNBlock(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:195
static constexpr __device__ auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:526
static constexpr auto I4
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:123
static constexpr __device__ auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:534
__host__ static constexpr __device__ auto MakeGemmMmaTileDescriptor(const TileDesc_K0_MN_K1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:201
static constexpr auto I5
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:124
static constexpr __device__ auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:542
static __host__ auto CalculateBK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:173
static constexpr auto BK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:132
static __device__ auto MakeAGridDescriptor_AK0_M_AK1(index_t M, index_t MPad, index_t K, index_t KPad, index_t StrideA, index_t AK0)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:215
__host__ static constexpr __device__ auto MakeAMmaTileDescriptor_M0_M1_M2_K(const ABlockDesc_AK0_M_AK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:381
static constexpr auto I3
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:122
static __host__ auto CalculateKPadded(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:151
__host__ static __device__ auto MakeCGridDescriptor_M_N(index_t M, index_t MPad, index_t N, index_t NPad, index_t StrideC)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:398
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:146
__host__ static constexpr __device__ auto MakeBMmaTileDescriptor_N0_N1_N2_K(const BBlockDesc_BK0_N_BK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:390
static constexpr auto I7
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:126
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:585
static __device__ auto MakeBGridDescriptor_BK0_N_BK1(index_t K, index_t KPad, index_t N, index_t NPad, index_t StrideB, index_t BK0)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:297
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:134
Definition: xdlops_gemm.hpp:1126
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
Definition: integral_constant.hpp:20
Definition: functional2.hpp:33
Definition: device_base.hpp:51
Definition: unary_element_wise_operation.hpp:334