20 template <
typename Gr
idwiseGemm,
bool HasMainKBlockLoop>
22 #if CK_USE_LAUNCH_BOUNDS
27 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
29 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
31 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
33 GridwiseGemm::template Run<HasMainKBlockLoop>(
34 karg.p_a_grid, karg.p_b_grid, karg.p_c_grid, p_shared, karg);
41 template <
typename GridwiseGemm,
45 bool HasMainKBlockLoop>
47 #if CK_USE_LAUNCH_BOUNDS
51 const FloatB* __restrict__ p_b_grid,
52 FloatC* __restrict__ p_c_grid,
53 typename GridwiseGemm::Problem problem)
55 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
57 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
59 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
61 GridwiseGemm::template Run<HasMainKBlockLoop>(
62 p_a_grid, p_b_grid, p_c_grid, p_shared, problem);
72 template <
typename ALayout,
77 typename FloatGemmAcc,
78 typename FloatCShuffle,
80 typename AElementwiseOperation,
81 typename BElementwiseOperation,
82 typename CElementwiseOperation,
96 typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
97 typename ABlockTransferThreadClusterArrangeOrder,
98 typename ABlockTransferSrcAccessOrder,
99 index_t ABlockTransferSrcVectorDim,
100 index_t ABlockTransferSrcScalarPerVector,
101 index_t ABlockTransferDstScalarPerVector_AK1,
102 bool AThreadTransferSrcResetCoordinateAfterRun,
104 typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
105 typename BBlockTransferThreadClusterArrangeOrder,
106 typename BBlockTransferSrcAccessOrder,
107 index_t BBlockTransferSrcVectorDim,
108 index_t BBlockTransferSrcScalarPerVector,
109 index_t BBlockTransferDstScalarPerVector_BK1,
110 bool BThreadTransferSrcResetCoordinateAfterRun,
112 index_t CShuffleMXdlPerWavePerShuffle,
113 index_t CShuffleNXdlPerWavePerShuffle,
114 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
115 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
118 typename ComputeTypeA = FloatC,
119 typename ComputeTypeB = ComputeTypeA>
163 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
164 GemmSpec == GemmSpecialization::MNKPadding ||
165 GemmSpec == GemmSpecialization::KPadding ||
166 GemmSpec == GemmSpecialization::NKPadding)
180 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
181 GemmSpec == GemmSpecialization::MNKPadding ||
182 GemmSpec == GemmSpecialization::KPadding ||
183 GemmSpec == GemmSpecialization::MKPadding)
206 const auto a_grid_desc_mraw_kraw = [&]() {
207 if constexpr(is_same_v<tensor_layout::gemm::RowMajor, ALayout>)
211 else if constexpr(is_same_v<tensor_layout::gemm::ColumnMajor, ALayout>)
219 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
220 GemmSpec == GemmSpecialization::MNKPadding)
223 const auto a_grid_desc_m_k =
237 return a_grid_desc_ak0_m_ak1;
239 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
240 GemmSpec == GemmSpecialization::MNPadding)
244 a_grid_desc_mraw_kraw,
250 return a_grid_desc_ak0_m_ak1;
252 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
253 GemmSpec == GemmSpecialization::NKPadding)
257 a_grid_desc_mraw_kraw,
269 return a_grid_desc_ak0_m_ak1;
275 a_grid_desc_mraw_kraw,
281 return a_grid_desc_ak0_m_ak1;
288 const auto b_grid_desc_nraw_kraw = [&]() {
301 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
302 GemmSpec == GemmSpecialization::MNKPadding)
305 const auto b_grid_desc_n_k =
319 return b_grid_desc_bk0_n_bk1;
321 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
322 GemmSpec == GemmSpecialization::MNPadding)
326 b_grid_desc_nraw_kraw,
332 return b_grid_desc_bk0_n_bk1;
334 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
335 GemmSpec == GemmSpecialization::MKPadding)
339 b_grid_desc_nraw_kraw,
351 return b_grid_desc_bk0_n_bk1;
357 b_grid_desc_nraw_kraw,
363 return b_grid_desc_bk0_n_bk1;
367 __host__ __device__
static auto
370 const auto c_grid_desc_mraw_nraw = [&]() {
383 if constexpr(GemmSpec == GemmSpecialization::MNPadding ||
384 GemmSpec == GemmSpecialization::MNKPadding)
393 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
394 GemmSpec == GemmSpecialization::MKPadding)
398 c_grid_desc_mraw_nraw,
403 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
404 GemmSpec == GemmSpecialization::NKPadding)
408 c_grid_desc_mraw_nraw,
416 return c_grid_desc_mraw_nraw;
446 std::cout <<
"problem {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
449 <<
"KP:" <<
KPadded <<
", " <<
"AK0:" <<
AK0 <<
", " <<
"BK0:" <<
BK0 <<
", "
450 <<
"MBlock: " <<
MBlock <<
", " <<
"NBlock: " <<
NBlock <<
"}" << std::endl;
472 const FloatB* p_b_grid_,
480 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
494 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
514 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
515 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
517 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
524 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock;
537 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
540 b_block_desc_bk0_n_bk1.GetElementSpaceSize(), max_lds_align);
543 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
546 constexpr
auto c_block_size =
547 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize();
549 return math::max((a_block_space_size_aligned *
sizeof(ComputeTypeA) +
550 b_block_space_size_aligned *
sizeof(ComputeTypeB)),
551 c_block_size *
sizeof(FloatCShuffle));
558 return ck::tensor_operation::device::IsValidGemmCompilationParameter<
567 CGlobalMemoryDataOperation>();
573 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
574 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
575 "Invalid tuning param!");
582 if(!(problem.
M % MPerBlock == 0))
593 if(!(problem.
N % NPerBlock == 0))
612 if(!(problem.
K % AK1Value == 0) || !(problem.
K % BK1Value == 0))
620 if(problem.
K % ABlockTransferSrcScalarPerVector != 0)
627 if(problem.
M % ABlockTransferSrcScalarPerVector != 0)
635 if(problem.
N % BBlockTransferSrcScalarPerVector != 0)
642 if(problem.
K % BBlockTransferSrcScalarPerVector != 0)
650 if(problem.
N % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
657 if(problem.
M % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
664 const auto num_k_loop = (
CalculateAK0(problem.
K) * AK1Value) / KPerBlock;
666 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
677 const index_t num_loop = K / KPerBlock;
679 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
682 template <
typename CGr
idDesc>
693 return c_grid_desc_mblock_mperblock_nblock_nperblock;
699 template <
bool HasMainKBlockLoop>
700 __device__
static void Run(
const FloatA* __restrict__ p_a_grid,
701 const FloatB* __restrict__ p_b_grid,
702 FloatC* __restrict__ p_c_grid,
703 void* __restrict__ p_shared,
713 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
717 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
718 p_a_grid, a_grid_desc_ak0_m_ak1.GetElementSpaceSize());
719 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
720 p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize());
721 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
722 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
724 const AElementwiseOperation a_element_op{};
725 const BElementwiseOperation b_element_op{};
726 const CElementwiseOperation c_element_op{};
731 const auto block_work_idx =
734 if(!block_2_ctile_map.ValidCTileIndex(
736 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
737 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
743 const index_t m_block_data_idx_on_grid =
744 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
746 const index_t n_block_data_idx_on_grid =
747 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
759 auto a_blockwise_copy =
761 AElementwiseOperation,
765 ABlockTransferThreadClusterLengths_AK0_M_AK1,
766 ABlockTransferThreadClusterArrangeOrder,
769 decltype(a_grid_desc_ak0_m_ak1),
770 decltype(a_block_desc_ak0_m_ak1),
771 ABlockTransferSrcAccessOrder,
773 ABlockTransferSrcVectorDim,
775 ABlockTransferSrcScalarPerVector,
776 ABlockTransferDstScalarPerVector_AK1,
779 AThreadTransferSrcResetCoordinateAfterRun,
781 NumGemmKPrefetchStage>(
782 a_grid_desc_ak0_m_ak1,
785 a_block_desc_ak0_m_ak1,
790 auto b_blockwise_copy =
792 BElementwiseOperation,
796 BBlockTransferThreadClusterLengths_BK0_N_BK1,
797 BBlockTransferThreadClusterArrangeOrder,
800 decltype(b_grid_desc_bk0_n_bk1),
801 decltype(b_block_desc_bk0_n_bk1),
802 BBlockTransferSrcAccessOrder,
804 BBlockTransferSrcVectorDim,
806 BBlockTransferSrcScalarPerVector,
807 BBlockTransferDstScalarPerVector_BK1,
810 BThreadTransferSrcResetCoordinateAfterRun,
812 NumGemmKPrefetchStage>(
813 b_grid_desc_bk0_n_bk1,
816 b_block_desc_bk0_n_bk1,
828 constexpr
bool is_single_rate_mfma =
836 constexpr
auto is_scale_mfma =
false;
843 is_scale_mfma>::selected_mfma.k_per_blk);
850 decltype(a_block_desc_ak0_m_ak1),
851 decltype(b_block_desc_bk0_n_bk1),
859 auto c_thread_buf = blockwise_gemm.GetCThreadBuffer();
863 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
865 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
866 static_cast<ComputeTypeA*
>(p_shared), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
868 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
869 static_cast<ComputeTypeB*
>(p_shared) + a_block_space_size_aligned,
870 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
876 static_assert(std::is_default_constructible_v<GridwiseGemmPipe>);
879 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
880 (a_grid_desc_ak0_m_ak1.GetLength(
I0) * a_grid_desc_ak0_m_ak1.GetLength(
I2)) /
883 gridwise_gemm_pipeline.template Run<HasMainKBlockLoop>(a_grid_desc_ak0_m_ak1,
884 a_block_desc_ak0_m_ak1,
888 a_block_slice_copy_step,
889 b_grid_desc_bk0_n_bk1,
890 b_block_desc_bk0_n_bk1,
894 b_block_slice_copy_step,
897 num_k_block_main_loop);
901 static_assert(MXdlPerWave % CShuffleMXdlPerWavePerShuffle == 0 &&
902 NXdlPerWave % CShuffleNXdlPerWavePerShuffle == 0,
905 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
906 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
909 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
910 blockwise_gemm.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
914 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp =
915 blockwise_gemm.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
917 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I0);
918 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I1);
919 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I2);
920 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I3);
921 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I4);
922 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I5);
923 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I6);
924 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I7);
926 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
929 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
930 static_cast<FloatCShuffle*
>(p_shared),
931 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
934 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
954 const auto c_thread_mtx_on_block =
955 blockwise_gemm.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
957 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
958 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
960 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
966 const auto m_thread_data_on_block_idx =
967 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
970 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
976 const auto n_thread_data_on_block_idx =
977 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
981 auto c_thread_copy_vgpr_to_lds =
984 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
985 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
987 Sequence<CShuffleMXdlPerWavePerShuffle,
988 CShuffleNXdlPerWavePerShuffle,
1001 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1004 m_thread_data_on_block_idx[
I1],
1005 n_thread_data_on_block_idx[
I1],
1006 m_thread_data_on_block_idx[
I2],
1007 m_thread_data_on_block_idx[
I3],
1008 m_thread_data_on_block_idx[
I4],
1009 n_thread_data_on_block_idx[
I2]),
1015 CElementwiseOperation,
1016 CGlobalMemoryDataOperation,
1018 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1020 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>,
1021 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1025 decltype(c_shuffle_block_desc_mblock_mperblock_nblock_nperblock),
1026 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1029 CShuffleBlockTransferScalarPerVector_NPerBlock,
1032 {c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1034 c_grid_desc_mblock_mperblock_nblock_nperblock,
1039 constexpr
auto sfc_c_vgpr =
1042 Sequence<CShuffleMXdlPerWavePerShuffle,
1043 CShuffleNXdlPerWavePerShuffle,
1052 constexpr
auto sfc_c_global =
1056 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1058 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{};
1060 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
1062 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1069 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1070 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1072 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1073 c_shuffle_block_buf);
1079 c_shuffle_block_copy_lds_to_global.Run(
1080 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1081 c_shuffle_block_buf,
1082 c_grid_desc_mblock_mperblock_nblock_nperblock,
1085 if constexpr(access_id < num_access - 1)
1087 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1090 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1091 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:31
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:30
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:620
__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:277
__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:151
constexpr detail::ignore_t ignore
Definition: ignore.hpp:20
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:47
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__host__ constexpr __device__ auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:211
remove_cv_t< remove_reference_t< T > > remove_cvref_t
Definition: type.hpp:297
__host__ constexpr __device__ auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:90
LoopScheduler
Definition: loop_scheduler.hpp:15
int32_t index_t
Definition: ck.hpp:299
__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:470
__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:471
const FloatB * p_b_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:488
const FloatA * p_a_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:487
FloatC * p_c_grid
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:489
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:421
index_t NPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:460
index_t N
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:454
index_t StrideA
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:456
index_t StrideB
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:457
index_t NBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:465
index_t BK0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:463
index_t KPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:461
index_t MPadded
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:459
index_t K
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:455
index_t MBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:464
__host__ void Print() const
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:444
__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:422
index_t M
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:453
index_t AK0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:462
index_t StrideC
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:458
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:121
static constexpr __device__ auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:504
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:571
static constexpr auto BK1Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:135
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:149
static __host__ auto CalculateAK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:159
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:144
static constexpr auto I6
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:128
static constexpr auto I1
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:123
static constexpr auto I2
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:124
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:675
static constexpr auto I5
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:127
__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:368
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:137
static __host__ auto CalculateBK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:176
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:700
static __host__ auto CalculateNBlock(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:198
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:285
static __device__ constexpr bool IsValidCompilationParameter()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:556
static constexpr auto BK0Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:133
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:203
static constexpr auto I7
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:129
static constexpr auto I3
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:125
static __host__ auto CalculateKPadded(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:154
static constexpr auto AK1Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:134
static __host__ auto CalculateMBlock(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:193
static constexpr auto AK0Number
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:132
static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:527
static constexpr auto I0
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:122
static constexpr __device__ auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:512
static constexpr auto I4
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:126
static constexpr __device__ auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:496
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:494
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:683
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v1.hpp:139
Selects the appropriate MFMA instruction type and configuration for given data types and tile sizes o...
Definition: xdlops_gemm.hpp:1208
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:197
Definition: unary_element_wise_operation.hpp:334