20 template <
typename Gr
idwiseGemm,
bool HasMainKBlockLoop, index_t TailNum = 3>
22 #if CK_USE_LAUNCH_BOUNDS
28 #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__) || \
32 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
33 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
35 GridwiseGemm::template Run<HasMainKBlockLoop, TailNum>(
36 karg.p_a_grid, karg.p_b_grid, karg.p_c_grid, p_shared_0, p_shared_1, karg);
42 template <
typename GridwiseGemm,
46 bool HasMainKBlockLoop>
48 #if CK_USE_LAUNCH_BOUNDS
52 const FloatB* p_b_grid,
54 typename GridwiseGemm::Problem problem)
56 #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__) || \
58 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
59 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
61 GridwiseGemm::template Run<HasMainKBlockLoop>(
62 p_a_grid, p_b_grid, p_c_grid, p_shared_0, p_shared_1, problem);
71 template <
typename ALayout,
76 typename FloatGemmAcc,
77 typename FloatCShuffle,
79 typename AElementwiseOperation,
80 typename BElementwiseOperation,
81 typename CElementwiseOperation,
95 typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
96 typename ABlockTransferThreadClusterArrangeOrder,
97 typename ABlockTransferSrcAccessOrder,
98 index_t ABlockTransferSrcVectorDim,
99 index_t ABlockTransferSrcScalarPerVector,
100 index_t ABlockTransferDstScalarPerVector_AK1,
101 bool AThreadTransferSrcResetCoordinateAfterRun,
103 typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
104 typename BBlockTransferThreadClusterArrangeOrder,
105 typename BBlockTransferSrcAccessOrder,
106 index_t BBlockTransferSrcVectorDim,
107 index_t BBlockTransferSrcScalarPerVector,
108 index_t BBlockTransferDstScalarPerVector_BK1,
109 bool BThreadTransferSrcResetCoordinateAfterRun,
111 index_t CShuffleMXdlPerWavePerShuffle,
112 index_t CShuffleNXdlPerWavePerShuffle,
113 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
114 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
117 typename ComputeTypeA = FloatC,
118 typename ComputeTypeB = ComputeTypeA>
162 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
163 GemmSpec == GemmSpecialization::MNKPadding ||
164 GemmSpec == GemmSpecialization::KPadding ||
165 GemmSpec == GemmSpecialization::NKPadding)
179 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
180 GemmSpec == GemmSpecialization::MNKPadding ||
181 GemmSpec == GemmSpecialization::KPadding ||
182 GemmSpec == GemmSpecialization::MKPadding)
202 template <index_t MNXdlPerWave, index_t MNWaves, index_t MNPerXdl,
typename TileDesc_K0_MN_K1>
220 const auto a_grid_desc_mraw_kraw = [&]() {
221 if constexpr(is_same_v<tensor_layout::gemm::RowMajor, ALayout>)
225 else if constexpr(is_same_v<tensor_layout::gemm::ColumnMajor, ALayout>)
233 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
234 GemmSpec == GemmSpecialization::MNKPadding)
237 const auto a_grid_desc_m_k =
251 return a_grid_desc_ak0_m_ak1;
253 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
254 GemmSpec == GemmSpecialization::MNPadding)
258 a_grid_desc_mraw_kraw,
264 return a_grid_desc_ak0_m_ak1;
266 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
267 GemmSpec == GemmSpecialization::NKPadding)
271 a_grid_desc_mraw_kraw,
283 return a_grid_desc_ak0_m_ak1;
289 a_grid_desc_mraw_kraw,
295 return a_grid_desc_ak0_m_ak1;
302 const auto b_grid_desc_nraw_kraw = [&]() {
315 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
316 GemmSpec == GemmSpecialization::MNKPadding)
319 const auto b_grid_desc_n_k =
333 return b_grid_desc_bk0_n_bk1;
335 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
336 GemmSpec == GemmSpecialization::MNPadding)
340 b_grid_desc_nraw_kraw,
346 return b_grid_desc_bk0_n_bk1;
348 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
349 GemmSpec == GemmSpecialization::MKPadding)
353 b_grid_desc_nraw_kraw,
365 return b_grid_desc_bk0_n_bk1;
371 b_grid_desc_nraw_kraw,
377 return b_grid_desc_bk0_n_bk1;
381 template <
typename ABlockDesc_AK0_M_AK1>
382 __host__ __device__
static constexpr
auto
385 constexpr
index_t MWaves = MPerBlock / (MXdlPerWave * MPerXdl);
387 return MakeGemmMmaTileDescriptor<MXdlPerWave, MWaves, MPerXdl>(ABlockDesc_AK0_M_AK1{});
390 template <
typename BBlockDesc_BK0_N_BK1>
391 __host__ __device__
static constexpr
auto
394 constexpr
index_t NWaves = NPerBlock / (NXdlPerWave * NPerXdl);
396 return MakeGemmMmaTileDescriptor<NXdlPerWave, NWaves, NPerXdl>(BBlockDesc_BK0_N_BK1{});
399 __host__ __device__
static auto
402 const auto c_grid_desc_mraw_nraw = [&]() {
415 if constexpr(GemmSpec == GemmSpecialization::MNPadding ||
416 GemmSpec == GemmSpecialization::MNKPadding)
425 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
426 GemmSpec == GemmSpecialization::MKPadding)
430 c_grid_desc_mraw_nraw,
435 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
436 GemmSpec == GemmSpecialization::NKPadding)
440 c_grid_desc_mraw_nraw,
448 return c_grid_desc_mraw_nraw;
478 std::cout <<
"problem {"
488 <<
"AK0:" <<
AK0 <<
", "
489 <<
"BK0:" <<
BK0 <<
", "
490 <<
"MBlock: " <<
MBlock <<
", "
491 <<
"NBlock: " <<
NBlock <<
"}" << std::endl;
513 const FloatB* p_b_grid_,
521 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
535 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
555 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
556 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
558 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
565 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock;
578 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
581 b_block_desc_bk0_n_bk1.GetElementSpaceSize(), max_lds_align);
584 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
587 constexpr
auto c_block_size =
588 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize();
590 return math::max((a_block_space_size_aligned *
sizeof(ComputeTypeA) +
591 b_block_space_size_aligned *
sizeof(ComputeTypeB)),
592 c_block_size *
sizeof(FloatCShuffle));
598 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
599 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
600 "Invalid tuning param!");
607 if(!(problem.
M % MPerBlock == 0))
618 if(!(problem.
N % NPerBlock == 0))
637 if(!(problem.
K % AK1Value == 0) || !(problem.
K % BK1Value == 0))
645 if(problem.
K % ABlockTransferSrcScalarPerVector != 0)
652 if(problem.
M % ABlockTransferSrcScalarPerVector != 0)
660 if(problem.
N % BBlockTransferSrcScalarPerVector != 0)
667 if(problem.
K % BBlockTransferSrcScalarPerVector != 0)
675 if(problem.
N % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
682 if(problem.
M % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
689 const auto num_k_loop = (
CalculateAK0(problem.
K) * AK1Value) / KPerBlock;
702 const index_t num_loop = K / KPerBlock;
709 const index_t num_loop = K / KPerBlock;
711 if(num_loop % 2 == 1)
717 template <
typename CGr
idDesc>
728 return c_grid_desc_mblock_mperblock_nblock_nperblock;
735 template <
bool HasMainKBlockLoop, index_t TailNum = 3>
736 __device__
static void Run(
const FloatA* p_a_grid,
737 const FloatB* p_b_grid,
750 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
754 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
755 p_a_grid, a_grid_desc_ak0_m_ak1.GetElementSpaceSize());
756 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
757 p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize());
758 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
759 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
761 const AElementwiseOperation a_element_op{};
762 const BElementwiseOperation b_element_op{};
763 const CElementwiseOperation c_element_op{};
768 const auto block_work_idx =
771 if(!block_2_ctile_map.ValidCTileIndex(
773 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
774 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
779 if(threadIdx.x == 0){
780 printf(
"Hardware assigned No. %03d workgroup of logical C tile (%02d, %02d) on %d th XCC Die, %d th SE, %d th CU\n",
790 const index_t m_block_data_idx_on_grid =
791 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
793 const index_t n_block_data_idx_on_grid =
794 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
806 auto a_blockwise_copy =
808 AElementwiseOperation,
812 ABlockTransferThreadClusterLengths_AK0_M_AK1,
813 ABlockTransferThreadClusterArrangeOrder,
816 decltype(a_grid_desc_ak0_m_ak1),
817 decltype(a_block_desc_ak0_m_ak1),
818 ABlockTransferSrcAccessOrder,
820 ABlockTransferSrcVectorDim,
822 ABlockTransferSrcScalarPerVector,
823 ABlockTransferDstScalarPerVector_AK1,
826 AThreadTransferSrcResetCoordinateAfterRun,
828 a_grid_desc_ak0_m_ak1,
831 a_block_desc_ak0_m_ak1,
836 auto b_blockwise_copy =
838 BElementwiseOperation,
842 BBlockTransferThreadClusterLengths_BK0_N_BK1,
843 BBlockTransferThreadClusterArrangeOrder,
846 decltype(b_grid_desc_bk0_n_bk1),
847 decltype(b_block_desc_bk0_n_bk1),
848 BBlockTransferSrcAccessOrder,
850 BBlockTransferSrcVectorDim,
852 BBlockTransferSrcScalarPerVector,
853 BBlockTransferDstScalarPerVector_BK1,
856 BThreadTransferSrcResetCoordinateAfterRun,
858 b_grid_desc_bk0_n_bk1,
861 b_block_desc_bk0_n_bk1,
892 decltype(a_block_desc_ak0_m_ak1),
893 decltype(b_block_desc_bk0_n_bk1),
905 auto c_thread_buf = blockwise_gemm_pipeline.GetCThreadBuffer();
909 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
911 auto a_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
912 static_cast<ComputeTypeA*
>(p_shared_0), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
914 auto b_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
915 static_cast<ComputeTypeB*
>(p_shared_0) + a_block_space_size_aligned,
916 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
918 auto a_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
919 static_cast<ComputeTypeA*
>(p_shared_1), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
921 auto b_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
922 static_cast<ComputeTypeB*
>(p_shared_1) + a_block_space_size_aligned,
923 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
925 auto a_block_bufs =
make_tuple(a_block_buf_ping, a_block_buf_pong);
926 auto b_block_bufs =
make_tuple(b_block_buf_ping, b_block_buf_pong);
932 static_assert(std::is_default_constructible_v<GridwiseGemmPipe>);
935 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
936 (a_grid_desc_ak0_m_ak1.GetLength(
I0) * a_grid_desc_ak0_m_ak1.GetLength(
I2)) /
939 blockwise_gemm_pipeline.template Run<HasMainKBlockLoop, TailNum>(a_grid_desc_ak0_m_ak1,
940 a_block_desc_ak0_m_ak1,
944 a_block_slice_copy_step,
945 b_grid_desc_bk0_n_bk1,
946 b_block_desc_bk0_n_bk1,
950 b_block_slice_copy_step,
952 num_k_block_main_loop);
956 static_assert(MXdlPerWave % CShuffleMXdlPerWavePerShuffle == 0 &&
957 NXdlPerWave % CShuffleNXdlPerWavePerShuffle == 0,
960 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
961 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
964 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
965 blockwise_gemm_pipeline.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
969 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp =
970 blockwise_gemm_pipeline.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
972 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I0);
973 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I1);
974 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I2);
975 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I3);
976 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I4);
977 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I5);
978 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I6);
979 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I7);
981 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
984 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
985 static_cast<FloatCShuffle*
>(p_shared_0),
986 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
989 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1009 const auto c_thread_mtx_on_block =
1010 blockwise_gemm_pipeline.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
1012 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
1013 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
1015 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
1021 const auto m_thread_data_on_block_idx =
1022 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
1025 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
1031 const auto n_thread_data_on_block_idx =
1032 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
1036 auto c_thread_copy_vgpr_to_lds =
1039 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1040 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1042 Sequence<CShuffleMXdlPerWavePerShuffle,
1043 CShuffleNXdlPerWavePerShuffle,
1056 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1059 m_thread_data_on_block_idx[
I1],
1060 n_thread_data_on_block_idx[
I1],
1061 m_thread_data_on_block_idx[
I2],
1062 m_thread_data_on_block_idx[
I3],
1063 m_thread_data_on_block_idx[
I4],
1064 n_thread_data_on_block_idx[
I2]),
1070 CElementwiseOperation,
1071 CGlobalMemoryDataOperation,
1073 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1075 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>,
1076 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1080 decltype(c_shuffle_block_desc_mblock_mperblock_nblock_nperblock),
1081 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1084 CShuffleBlockTransferScalarPerVector_NPerBlock,
1087 {c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1089 c_grid_desc_mblock_mperblock_nblock_nperblock,
1094 constexpr
auto sfc_c_vgpr =
1097 Sequence<CShuffleMXdlPerWavePerShuffle,
1098 CShuffleNXdlPerWavePerShuffle,
1107 constexpr
auto sfc_c_global =
1111 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1113 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{};
1115 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
1117 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1124 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1125 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1127 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1128 c_shuffle_block_buf);
1134 c_shuffle_block_copy_lds_to_global.Run(
1135 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1136 c_shuffle_block_buf,
1137 c_grid_desc_mblock_mperblock_nblock_nperblock,
1140 if constexpr(access_id < num_access - 1)
1142 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1145 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1146 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:33
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:267
__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:429
__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:22
__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:300
__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:289
__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:17
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:281
Definition: blockwise_gemm_pipeline_xdlops.hpp:103
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:511
const FloatB * p_b_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:529
__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:512
FloatC * p_c_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:530
const FloatA * p_a_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:528
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:453
index_t StrideB
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:498
index_t StrideC
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:499
index_t KPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:502
index_t N
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:495
index_t M
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:494
index_t MBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:505
index_t NBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:506
__host__ void Print() const
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:476
index_t BK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:504
index_t MPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:500
index_t K
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:496
index_t StrideA
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:497
index_t AK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:503
__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:454
index_t NPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:501
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:120
static constexpr auto I6
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:127
static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:568
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:143
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:138
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:700
static constexpr __host__ index_t CalculateKBlockLoopTailNum(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:707
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:535
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:718
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:736
static constexpr auto AK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:133
static constexpr auto I1
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:122
static constexpr auto I2
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:123
static constexpr auto AK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:131
static constexpr auto I0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:121
static __host__ auto CalculateMBlock(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:192
static constexpr auto BK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:132
static __host__ auto CalculateAK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:158
static __host__ auto CalculateNBlock(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:197
static constexpr __device__ auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:537
static constexpr auto I4
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:125
static constexpr __device__ auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:545
__host__ static constexpr __device__ auto MakeGemmMmaTileDescriptor(const TileDesc_K0_MN_K1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:203
static constexpr auto I5
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:126
static constexpr __device__ auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:553
static __host__ auto CalculateBK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:175
static constexpr auto BK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:134
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:217
__host__ static constexpr __device__ auto MakeAMmaTileDescriptor_M0_M1_M2_K(const ABlockDesc_AK0_M_AK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:383
static constexpr auto I3
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:124
static __host__ auto CalculateKPadded(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:153
__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:400
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:148
__host__ static constexpr __device__ auto MakeBMmaTileDescriptor_N0_N1_N2_K(const BBlockDesc_BK0_N_BK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:392
static constexpr auto I7
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:128
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:596
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:299
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:136
Definition: xdlops_gemm.hpp:886
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:10
Definition: functional2.hpp:31
Definition: device_base.hpp:50
Definition: unary_element_wise_operation.hpp:241