20 template <
typename Gr
idwiseGemm,
bool HasMainKBlockLoop>
22 #if CK_USE_LAUNCH_BOUNDS
27 #if(defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__))
28 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
30 GridwiseGemm::template Run<HasMainKBlockLoop>(
31 karg.p_a_grid, karg.p_b_grid, karg.p_c_grid, p_shared, karg);
37 template <
typename GridwiseGemm,
41 bool HasMainKBlockLoop>
43 #if CK_USE_LAUNCH_BOUNDS
47 const FloatB* __restrict__ p_b_grid,
48 FloatC* __restrict__ p_c_grid,
49 typename GridwiseGemm::Problem problem)
51 #if(defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__))
52 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
54 GridwiseGemm::template Run<HasMainKBlockLoop>(p_a_grid, p_b_grid, p_c_grid, p_shared, problem);
63 template <
typename ALayout,
68 typename FloatGemmAcc,
69 typename FloatCShuffle,
71 typename AElementwiseOperation,
72 typename BElementwiseOperation,
73 typename CElementwiseOperation,
87 typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
88 typename ABlockTransferThreadClusterArrangeOrder,
89 typename ABlockTransferSrcAccessOrder,
90 index_t ABlockTransferSrcVectorDim,
91 index_t ABlockTransferSrcScalarPerVector,
92 index_t ABlockTransferDstScalarPerVector_AK1,
93 bool AThreadTransferSrcResetCoordinateAfterRun,
95 typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
96 typename BBlockTransferThreadClusterArrangeOrder,
97 typename BBlockTransferSrcAccessOrder,
98 index_t BBlockTransferSrcVectorDim,
99 index_t BBlockTransferSrcScalarPerVector,
100 index_t BBlockTransferDstScalarPerVector_BK1,
101 bool BThreadTransferSrcResetCoordinateAfterRun,
103 index_t CShuffleMXdlPerWavePerShuffle,
104 index_t CShuffleNXdlPerWavePerShuffle,
105 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
106 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
109 typename ComputeTypeA = FloatC,
110 typename ComputeTypeB = ComputeTypeA>
154 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
155 GemmSpec == GemmSpecialization::MNKPadding ||
156 GemmSpec == GemmSpecialization::KPadding ||
157 GemmSpec == GemmSpecialization::NKPadding)
171 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
172 GemmSpec == GemmSpecialization::MNKPadding ||
173 GemmSpec == GemmSpecialization::KPadding ||
174 GemmSpec == GemmSpecialization::MKPadding)
197 const auto a_grid_desc_mraw_kraw = [&]() {
198 if constexpr(is_same_v<tensor_layout::gemm::RowMajor, ALayout>)
202 else if constexpr(is_same_v<tensor_layout::gemm::ColumnMajor, ALayout>)
210 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
211 GemmSpec == GemmSpecialization::MNKPadding)
214 const auto a_grid_desc_m_k =
228 return a_grid_desc_ak0_m_ak1;
230 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
231 GemmSpec == GemmSpecialization::MNPadding)
235 a_grid_desc_mraw_kraw,
241 return a_grid_desc_ak0_m_ak1;
243 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
244 GemmSpec == GemmSpecialization::NKPadding)
248 a_grid_desc_mraw_kraw,
260 return a_grid_desc_ak0_m_ak1;
266 a_grid_desc_mraw_kraw,
272 return a_grid_desc_ak0_m_ak1;
279 const auto b_grid_desc_nraw_kraw = [&]() {
292 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
293 GemmSpec == GemmSpecialization::MNKPadding)
296 const auto b_grid_desc_n_k =
310 return b_grid_desc_bk0_n_bk1;
312 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
313 GemmSpec == GemmSpecialization::MNPadding)
317 b_grid_desc_nraw_kraw,
323 return b_grid_desc_bk0_n_bk1;
325 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
326 GemmSpec == GemmSpecialization::MKPadding)
330 b_grid_desc_nraw_kraw,
342 return b_grid_desc_bk0_n_bk1;
348 b_grid_desc_nraw_kraw,
354 return b_grid_desc_bk0_n_bk1;
358 __host__ __device__
static auto
361 const auto c_grid_desc_mraw_nraw = [&]() {
374 if constexpr(GemmSpec == GemmSpecialization::MNPadding ||
375 GemmSpec == GemmSpecialization::MNKPadding)
384 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
385 GemmSpec == GemmSpecialization::MKPadding)
389 c_grid_desc_mraw_nraw,
394 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
395 GemmSpec == GemmSpecialization::NKPadding)
399 c_grid_desc_mraw_nraw,
407 return c_grid_desc_mraw_nraw;
437 std::cout <<
"problem {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
440 <<
"KP:" <<
KPadded <<
", " <<
"AK0:" <<
AK0 <<
", " <<
"BK0:" <<
BK0 <<
", "
441 <<
"MBlock: " <<
MBlock <<
", " <<
"NBlock: " <<
NBlock <<
"}" << std::endl;
463 const FloatB* p_b_grid_,
471 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
485 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
505 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
506 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
508 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
515 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock;
528 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
531 b_block_desc_bk0_n_bk1.GetElementSpaceSize(), max_lds_align);
534 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
537 constexpr
auto c_block_size =
538 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize();
540 return math::max((a_block_space_size_aligned *
sizeof(ComputeTypeA) +
541 b_block_space_size_aligned *
sizeof(ComputeTypeB)),
542 c_block_size *
sizeof(FloatCShuffle));
548 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
549 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
550 "Invalid tuning param!");
557 if(!(problem.
M % MPerBlock == 0))
568 if(!(problem.
N % NPerBlock == 0))
587 if(!(problem.
K % AK1Value == 0) || !(problem.
K % BK1Value == 0))
595 if(problem.
K % ABlockTransferSrcScalarPerVector != 0)
602 if(problem.
M % ABlockTransferSrcScalarPerVector != 0)
610 if(problem.
N % BBlockTransferSrcScalarPerVector != 0)
617 if(problem.
K % BBlockTransferSrcScalarPerVector != 0)
625 if(problem.
N % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
632 if(problem.
M % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
639 const auto num_k_loop = (
CalculateAK0(problem.
K) * AK1Value) / KPerBlock;
641 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
652 const index_t num_loop = K / KPerBlock;
654 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
657 template <
typename CGr
idDesc>
668 return c_grid_desc_mblock_mperblock_nblock_nperblock;
674 template <
bool HasMainKBlockLoop>
675 __device__
static void Run(
const FloatA* __restrict__ p_a_grid,
676 const FloatB* __restrict__ p_b_grid,
677 FloatC* __restrict__ p_c_grid,
678 void* __restrict__ p_shared,
688 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
692 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
693 p_a_grid, a_grid_desc_ak0_m_ak1.GetElementSpaceSize());
694 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
695 p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize());
696 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
697 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
699 const AElementwiseOperation a_element_op{};
700 const BElementwiseOperation b_element_op{};
701 const CElementwiseOperation c_element_op{};
706 const auto block_work_idx =
709 if(!block_2_ctile_map.ValidCTileIndex(
711 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
712 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
718 const index_t m_block_data_idx_on_grid =
719 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
721 const index_t n_block_data_idx_on_grid =
722 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
734 auto a_blockwise_copy =
736 AElementwiseOperation,
740 ABlockTransferThreadClusterLengths_AK0_M_AK1,
741 ABlockTransferThreadClusterArrangeOrder,
744 decltype(a_grid_desc_ak0_m_ak1),
745 decltype(a_block_desc_ak0_m_ak1),
746 ABlockTransferSrcAccessOrder,
748 ABlockTransferSrcVectorDim,
750 ABlockTransferSrcScalarPerVector,
751 ABlockTransferDstScalarPerVector_AK1,
754 AThreadTransferSrcResetCoordinateAfterRun,
756 NumGemmKPrefetchStage>(
757 a_grid_desc_ak0_m_ak1,
760 a_block_desc_ak0_m_ak1,
765 auto b_blockwise_copy =
767 BElementwiseOperation,
771 BBlockTransferThreadClusterLengths_BK0_N_BK1,
772 BBlockTransferThreadClusterArrangeOrder,
775 decltype(b_grid_desc_bk0_n_bk1),
776 decltype(b_block_desc_bk0_n_bk1),
777 BBlockTransferSrcAccessOrder,
779 BBlockTransferSrcVectorDim,
781 BBlockTransferSrcScalarPerVector,
782 BBlockTransferDstScalarPerVector_BK1,
785 BThreadTransferSrcResetCoordinateAfterRun,
787 NumGemmKPrefetchStage>(
788 b_grid_desc_bk0_n_bk1,
791 b_block_desc_bk0_n_bk1,
803 constexpr
bool is_single_rate_mfma =
811 constexpr
auto is_scale_mfma =
false;
818 is_scale_mfma>::selected_mfma.k_per_blk);
825 decltype(a_block_desc_ak0_m_ak1),
826 decltype(b_block_desc_bk0_n_bk1),
834 auto c_thread_buf = blockwise_gemm.GetCThreadBuffer();
838 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
840 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
841 static_cast<ComputeTypeA*
>(p_shared), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
843 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
844 static_cast<ComputeTypeB*
>(p_shared) + a_block_space_size_aligned,
845 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
851 static_assert(std::is_default_constructible_v<GridwiseGemmPipe>);
854 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
855 (a_grid_desc_ak0_m_ak1.GetLength(
I0) * a_grid_desc_ak0_m_ak1.GetLength(
I2)) /
858 gridwise_gemm_pipeline.template Run<HasMainKBlockLoop>(a_grid_desc_ak0_m_ak1,
859 a_block_desc_ak0_m_ak1,
863 a_block_slice_copy_step,
864 b_grid_desc_bk0_n_bk1,
865 b_block_desc_bk0_n_bk1,
869 b_block_slice_copy_step,
872 num_k_block_main_loop);
876 static_assert(MXdlPerWave % CShuffleMXdlPerWavePerShuffle == 0 &&
877 NXdlPerWave % CShuffleNXdlPerWavePerShuffle == 0,
880 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
881 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
884 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
885 blockwise_gemm.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
889 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp =
890 blockwise_gemm.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
892 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I0);
893 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I1);
894 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I2);
895 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I3);
896 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I4);
897 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I5);
898 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I6);
899 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I7);
901 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
904 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
905 static_cast<FloatCShuffle*
>(p_shared),
906 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
909 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
929 const auto c_thread_mtx_on_block =
930 blockwise_gemm.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
932 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
933 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
935 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
941 const auto m_thread_data_on_block_idx =
942 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
945 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
951 const auto n_thread_data_on_block_idx =
952 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
956 auto c_thread_copy_vgpr_to_lds =
959 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
960 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
962 Sequence<CShuffleMXdlPerWavePerShuffle,
963 CShuffleNXdlPerWavePerShuffle,
976 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
979 m_thread_data_on_block_idx[
I1],
980 n_thread_data_on_block_idx[
I1],
981 m_thread_data_on_block_idx[
I2],
982 m_thread_data_on_block_idx[
I3],
983 m_thread_data_on_block_idx[
I4],
984 n_thread_data_on_block_idx[
I2]),
990 CElementwiseOperation,
991 CGlobalMemoryDataOperation,
993 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
995 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>,
996 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1000 decltype(c_shuffle_block_desc_mblock_mperblock_nblock_nperblock),
1001 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1004 CShuffleBlockTransferScalarPerVector_NPerBlock,
1007 {c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1009 c_grid_desc_mblock_mperblock_nblock_nperblock,
1014 constexpr
auto sfc_c_vgpr =
1017 Sequence<CShuffleMXdlPerWavePerShuffle,
1018 CShuffleNXdlPerWavePerShuffle,
1027 constexpr
auto sfc_c_global =
1031 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1033 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{};
1035 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
1037 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1044 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1045 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1047 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1048 c_shuffle_block_buf);
1054 c_shuffle_block_copy_lds_to_global.Run(
1055 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1056 c_shuffle_block_buf,
1057 c_grid_desc_mblock_mperblock_nblock_nperblock,
1060 if constexpr(access_id < num_access - 1)
1062 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1065 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1066 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:30
#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
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
InMemoryDataOperationEnum
Definition: ck.hpp:276
__global__ void kernel_gemm_xdl_cshuffle_v1(typename GridwiseGemm::Argument karg)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:25
__host__ constexpr __device__ auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition: tensor_descriptor_helper.hpp:101
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__host__ constexpr __device__ auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:425
__host__ constexpr __device__ auto make_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
__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: gridwise_gemm_xdl_cshuffle_v1.hpp:461
__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_v1.hpp:462
const FloatB * p_b_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:479
const FloatA * p_a_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:478
FloatC * p_c_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:480
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:412
index_t NPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:451
index_t N
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:445
index_t StrideA
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:447
index_t StrideB
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:448
index_t NBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:456
index_t BK0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:454
index_t KPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:452
index_t MPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:450
index_t K
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:446
index_t MBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:455
__host__ void Print() const
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:435
__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_v1.hpp:413
index_t M
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:444
index_t AK0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:453
index_t StrideC
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:449
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:112
static constexpr __device__ auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:495
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:546
static constexpr auto BK1Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:126
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:140
static __host__ auto CalculateAK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:150
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:135
static constexpr auto I6
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:119
static constexpr auto I1
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:114
static constexpr auto I2
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:115
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:650
static constexpr auto I5
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:118
__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_v1.hpp:359
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:128
static __host__ auto CalculateBK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:167
static __device__ void Run(const FloatA *__restrict__ p_a_grid, const FloatB *__restrict__ p_b_grid, FloatC *__restrict__ p_c_grid, void *__restrict__ p_shared, const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:675
static __host__ auto CalculateNBlock(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:189
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_v1.hpp:276
static constexpr auto BK0Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:124
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_v1.hpp:194
static constexpr auto I7
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:120
static constexpr auto I3
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:116
static __host__ auto CalculateKPadded(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:145
static constexpr auto AK1Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:125
static __host__ auto CalculateMBlock(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:184
static constexpr auto AK0Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:123
static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:518
static constexpr auto I0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:113
static constexpr __device__ auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:503
static constexpr auto I4
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:117
static constexpr __device__ auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:487
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:485
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_v1.hpp:658
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:130
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