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__) || defined(__gfx11__) || \
30 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
34 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
35 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
37 GridwiseGemm::template Run<HasMainKBlockLoop, TailNum>(
38 karg.p_a_grid, karg.p_b_grid, karg.p_c_grid, p_shared_0, p_shared_1, karg);
45 template <
typename GridwiseGemm,
49 bool HasMainKBlockLoop>
51 #if CK_USE_LAUNCH_BOUNDS
55 const FloatB* p_b_grid,
57 typename GridwiseGemm::Problem problem)
59 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
61 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
63 __shared__
char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];
64 __shared__
char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];
66 GridwiseGemm::template Run<HasMainKBlockLoop>(
67 p_a_grid, p_b_grid, p_c_grid, p_shared_0, p_shared_1, problem);
77 template <
typename ALayout,
82 typename FloatGemmAcc,
83 typename FloatCShuffle,
85 typename AElementwiseOperation,
86 typename BElementwiseOperation,
87 typename CElementwiseOperation,
101 typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
102 typename ABlockTransferThreadClusterArrangeOrder,
103 typename ABlockTransferSrcAccessOrder,
104 index_t ABlockTransferSrcVectorDim,
105 index_t ABlockTransferSrcScalarPerVector,
106 index_t ABlockTransferDstScalarPerVector_AK1,
107 bool AThreadTransferSrcResetCoordinateAfterRun,
109 typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
110 typename BBlockTransferThreadClusterArrangeOrder,
111 typename BBlockTransferSrcAccessOrder,
112 index_t BBlockTransferSrcVectorDim,
113 index_t BBlockTransferSrcScalarPerVector,
114 index_t BBlockTransferDstScalarPerVector_BK1,
115 bool BThreadTransferSrcResetCoordinateAfterRun,
117 index_t CShuffleMXdlPerWavePerShuffle,
118 index_t CShuffleNXdlPerWavePerShuffle,
119 typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
120 index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
123 typename ComputeTypeA = FloatC,
124 typename ComputeTypeB = ComputeTypeA>
168 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
169 GemmSpec == GemmSpecialization::MNKPadding ||
170 GemmSpec == GemmSpecialization::KPadding ||
171 GemmSpec == GemmSpecialization::NKPadding)
185 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
186 GemmSpec == GemmSpecialization::MNKPadding ||
187 GemmSpec == GemmSpecialization::KPadding ||
188 GemmSpec == GemmSpecialization::MKPadding)
208 template <index_t MNXdlPerWave, index_t MNWaves, index_t MNPerXdl,
typename TileDesc_K0_MN_K1>
226 const auto a_grid_desc_mraw_kraw = [&]() {
227 if constexpr(is_same_v<tensor_layout::gemm::RowMajor, ALayout>)
231 else if constexpr(is_same_v<tensor_layout::gemm::ColumnMajor, ALayout>)
239 if constexpr(GemmSpec == GemmSpecialization::MKPadding ||
240 GemmSpec == GemmSpecialization::MNKPadding)
243 const auto a_grid_desc_m_k =
257 return a_grid_desc_ak0_m_ak1;
259 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
260 GemmSpec == GemmSpecialization::MNPadding)
264 a_grid_desc_mraw_kraw,
270 return a_grid_desc_ak0_m_ak1;
272 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
273 GemmSpec == GemmSpecialization::NKPadding)
277 a_grid_desc_mraw_kraw,
289 return a_grid_desc_ak0_m_ak1;
295 a_grid_desc_mraw_kraw,
301 return a_grid_desc_ak0_m_ak1;
308 const auto b_grid_desc_nraw_kraw = [&]() {
321 if constexpr(GemmSpec == GemmSpecialization::NKPadding ||
322 GemmSpec == GemmSpecialization::MNKPadding)
325 const auto b_grid_desc_n_k =
339 return b_grid_desc_bk0_n_bk1;
341 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
342 GemmSpec == GemmSpecialization::MNPadding)
346 b_grid_desc_nraw_kraw,
352 return b_grid_desc_bk0_n_bk1;
354 else if constexpr(GemmSpec == GemmSpecialization::KPadding ||
355 GemmSpec == GemmSpecialization::MKPadding)
359 b_grid_desc_nraw_kraw,
371 return b_grid_desc_bk0_n_bk1;
377 b_grid_desc_nraw_kraw,
383 return b_grid_desc_bk0_n_bk1;
387 template <
typename ABlockDesc_AK0_M_AK1>
388 __host__ __device__
static constexpr
auto
391 constexpr
index_t MWaves = MPerBlock / (MXdlPerWave * MPerXdl);
393 return MakeGemmMmaTileDescriptor<MXdlPerWave, MWaves, MPerXdl>(ABlockDesc_AK0_M_AK1{});
396 template <
typename BBlockDesc_BK0_N_BK1>
397 __host__ __device__
static constexpr
auto
400 constexpr
index_t NWaves = NPerBlock / (NXdlPerWave * NPerXdl);
402 return MakeGemmMmaTileDescriptor<NXdlPerWave, NWaves, NPerXdl>(BBlockDesc_BK0_N_BK1{});
405 __host__ __device__
static auto
408 const auto c_grid_desc_mraw_nraw = [&]() {
421 if constexpr(GemmSpec == GemmSpecialization::MNPadding ||
422 GemmSpec == GemmSpecialization::MNKPadding)
431 else if constexpr(GemmSpec == GemmSpecialization::MPadding ||
432 GemmSpec == GemmSpecialization::MKPadding)
436 c_grid_desc_mraw_nraw,
441 else if constexpr(GemmSpec == GemmSpecialization::NPadding ||
442 GemmSpec == GemmSpecialization::NKPadding)
446 c_grid_desc_mraw_nraw,
454 return c_grid_desc_mraw_nraw;
484 std::cout <<
"problem {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
487 <<
"KP:" <<
KPadded <<
", " <<
"AK0:" <<
AK0 <<
", " <<
"BK0:" <<
BK0 <<
", "
488 <<
"MBlock: " <<
MBlock <<
", " <<
"NBlock: " <<
NBlock <<
"}" << std::endl;
510 const FloatB* p_b_grid_,
518 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
532 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
552 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
553 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
555 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
562 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock;
575 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
578 b_block_desc_bk0_n_bk1.GetElementSpaceSize(), max_lds_align);
581 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
584 constexpr
auto c_block_size =
585 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize();
587 return math::max((a_block_space_size_aligned *
sizeof(ComputeTypeA) +
588 b_block_space_size_aligned *
sizeof(ComputeTypeB)),
589 c_block_size *
sizeof(FloatCShuffle));
596 return ck::tensor_operation::device::IsValidGemmCompilationParameter<
605 CGlobalMemoryDataOperation>();
611 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
612 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
613 "Invalid tuning param!");
620 if(!(problem.
M % MPerBlock == 0))
631 if(!(problem.
N % NPerBlock == 0))
650 if(!(problem.
K % AK1Value == 0) || !(problem.
K % BK1Value == 0))
658 if(problem.
K % ABlockTransferSrcScalarPerVector != 0)
665 if(problem.
M % ABlockTransferSrcScalarPerVector != 0)
673 if(problem.
N % BBlockTransferSrcScalarPerVector != 0)
680 if(problem.
K % BBlockTransferSrcScalarPerVector != 0)
688 if(problem.
N % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
695 if(problem.
M % CShuffleBlockTransferScalarPerVector_NPerBlock != 0)
702 const auto num_k_loop = (
CalculateAK0(problem.
K) * AK1Value) / KPerBlock;
715 const index_t num_loop = K / KPerBlock;
722 const index_t num_loop = K / KPerBlock;
724 if(num_loop % 2 == 1)
730 template <
typename CGr
idDesc>
741 return c_grid_desc_mblock_mperblock_nblock_nperblock;
748 template <
bool HasMainKBlockLoop, index_t TailNum = 3>
749 __device__
static void Run(
const FloatA* p_a_grid,
750 const FloatB* p_b_grid,
763 const auto c_grid_desc_mblock_mperblock_nblock_nperblock =
767 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
768 p_a_grid, a_grid_desc_ak0_m_ak1.GetElementSpaceSize());
769 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
770 p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize());
771 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
772 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
774 const AElementwiseOperation a_element_op{};
775 const BElementwiseOperation b_element_op{};
776 const CElementwiseOperation c_element_op{};
781 const auto block_work_idx =
784 if(!block_2_ctile_map.ValidCTileIndex(
786 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I0),
787 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(
I2))))
792 if(threadIdx.x == 0){
793 printf(
"Hardware assigned No. %03d workgroup of logical C tile (%02d, %02d) on %d th XCC Die, %d th SE, %d th CU\n",
803 const index_t m_block_data_idx_on_grid =
804 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
806 const index_t n_block_data_idx_on_grid =
807 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
819 auto a_blockwise_copy =
821 AElementwiseOperation,
825 ABlockTransferThreadClusterLengths_AK0_M_AK1,
826 ABlockTransferThreadClusterArrangeOrder,
829 decltype(a_grid_desc_ak0_m_ak1),
830 decltype(a_block_desc_ak0_m_ak1),
831 ABlockTransferSrcAccessOrder,
833 ABlockTransferSrcVectorDim,
835 ABlockTransferSrcScalarPerVector,
836 ABlockTransferDstScalarPerVector_AK1,
839 AThreadTransferSrcResetCoordinateAfterRun,
841 a_grid_desc_ak0_m_ak1,
844 a_block_desc_ak0_m_ak1,
849 auto b_blockwise_copy =
851 BElementwiseOperation,
855 BBlockTransferThreadClusterLengths_BK0_N_BK1,
856 BBlockTransferThreadClusterArrangeOrder,
859 decltype(b_grid_desc_bk0_n_bk1),
860 decltype(b_block_desc_bk0_n_bk1),
861 BBlockTransferSrcAccessOrder,
863 BBlockTransferSrcVectorDim,
865 BBlockTransferSrcScalarPerVector,
866 BBlockTransferDstScalarPerVector_BK1,
869 BThreadTransferSrcResetCoordinateAfterRun,
871 b_grid_desc_bk0_n_bk1,
874 b_block_desc_bk0_n_bk1,
886 constexpr
bool is_single_rate_mfma =
894 constexpr
auto is_scale_mfma =
false;
901 is_scale_mfma>::selected_mfma.k_per_blk);
919 decltype(a_block_desc_ak0_m_ak1),
920 decltype(b_block_desc_bk0_n_bk1),
932 auto c_thread_buf = blockwise_gemm_pipeline.GetCThreadBuffer();
936 a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align);
938 auto a_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
939 static_cast<ComputeTypeA*
>(p_shared_0), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
941 auto b_block_buf_ping = make_dynamic_buffer<AddressSpaceEnum::Lds>(
942 static_cast<ComputeTypeB*
>(p_shared_0) + a_block_space_size_aligned,
943 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
945 auto a_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
946 static_cast<ComputeTypeA*
>(p_shared_1), a_block_desc_ak0_m_ak1.GetElementSpaceSize());
948 auto b_block_buf_pong = make_dynamic_buffer<AddressSpaceEnum::Lds>(
949 static_cast<ComputeTypeB*
>(p_shared_1) + a_block_space_size_aligned,
950 b_block_desc_bk0_n_bk1.GetElementSpaceSize());
952 auto a_block_bufs =
make_tuple(a_block_buf_ping, a_block_buf_pong);
953 auto b_block_bufs =
make_tuple(b_block_buf_ping, b_block_buf_pong);
959 static_assert(std::is_default_constructible_v<GridwiseGemmPipe>);
962 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(
963 (a_grid_desc_ak0_m_ak1.GetLength(
I0) * a_grid_desc_ak0_m_ak1.GetLength(
I2)) /
966 blockwise_gemm_pipeline.template Run<HasMainKBlockLoop, TailNum>(a_grid_desc_ak0_m_ak1,
967 a_block_desc_ak0_m_ak1,
971 a_block_slice_copy_step,
972 b_grid_desc_bk0_n_bk1,
973 b_block_desc_bk0_n_bk1,
977 b_block_slice_copy_step,
979 num_k_block_main_loop);
983 static_assert(MXdlPerWave % CShuffleMXdlPerWavePerShuffle == 0 &&
984 NXdlPerWave % CShuffleNXdlPerWavePerShuffle == 0,
987 constexpr
index_t MWave = MPerBlock / (MXdlPerWave * MPerXdl);
988 constexpr
index_t NWave = NPerBlock / (NXdlPerWave * NPerXdl);
991 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
992 blockwise_gemm_pipeline.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
996 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp =
997 blockwise_gemm_pipeline.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
999 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I0);
1000 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I1);
1001 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I2);
1002 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I3);
1003 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I4);
1004 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I5);
1005 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I6);
1006 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2_tmp.GetLength(
I7);
1008 constexpr
auto c_shuffle_block_desc_mblock_mperblock_nblock_nperblock =
1011 auto c_shuffle_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
1012 static_cast<FloatCShuffle*
>(p_shared_0),
1013 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
1016 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1036 const auto c_thread_mtx_on_block =
1037 blockwise_gemm_pipeline.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
1039 const index_t m_thread_data_on_block = c_thread_mtx_on_block[
I0];
1040 const index_t n_thread_data_on_block = c_thread_mtx_on_block[
I1];
1042 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
1048 const auto m_thread_data_on_block_idx =
1049 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
1052 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
1058 const auto n_thread_data_on_block_idx =
1059 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
1063 auto c_thread_copy_vgpr_to_lds =
1066 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1067 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
1069 Sequence<CShuffleMXdlPerWavePerShuffle,
1070 CShuffleNXdlPerWavePerShuffle,
1083 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1086 m_thread_data_on_block_idx[
I1],
1087 n_thread_data_on_block_idx[
I1],
1088 m_thread_data_on_block_idx[
I2],
1089 m_thread_data_on_block_idx[
I3],
1090 m_thread_data_on_block_idx[
I4],
1091 n_thread_data_on_block_idx[
I2]),
1097 CElementwiseOperation,
1098 CGlobalMemoryDataOperation,
1100 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1102 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>,
1103 CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
1107 decltype(c_shuffle_block_desc_mblock_mperblock_nblock_nperblock),
1108 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
1111 CShuffleBlockTransferScalarPerVector_NPerBlock,
1114 {c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1116 c_grid_desc_mblock_mperblock_nblock_nperblock,
1121 constexpr
auto sfc_c_vgpr =
1124 Sequence<CShuffleMXdlPerWavePerShuffle,
1125 CShuffleNXdlPerWavePerShuffle,
1134 constexpr
auto sfc_c_global =
1138 CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl,
1140 CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{};
1142 constexpr
index_t num_access = sfc_c_vgpr.GetNumOfAccess();
1144 static_assert(num_access == sfc_c_global.GetNumOfAccess(),
"wrong!");
1151 c_thread_copy_vgpr_to_lds.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1152 sfc_c_vgpr.GetIndexTupleOfNumber(access_id),
1154 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
1155 c_shuffle_block_buf);
1161 c_shuffle_block_copy_lds_to_global.Run(
1162 c_shuffle_block_desc_mblock_mperblock_nblock_nperblock,
1163 c_shuffle_block_buf,
1164 c_grid_desc_mblock_mperblock_nblock_nperblock,
1167 if constexpr(access_id < num_access - 1)
1169 constexpr
auto c_global_step = sfc_c_global.GetForwardStep(access_id);
1172 c_shuffle_block_copy_lds_to_global.MoveDstSliceWindow(
1173 c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step);
#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
__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
__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: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
__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: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: block_to_ctile_map.hpp:271
__host__ static constexpr __device__ index_t CalculateGridSize(index_t M, index_t N)
Definition: block_to_ctile_map.hpp:283
Definition: blockwise_gemm_pipeline_xdlops.hpp:103
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:508
const FloatB * p_b_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:526
__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:509
FloatC * p_c_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:527
const FloatA * p_a_grid
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:525
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:459
index_t StrideB
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:495
index_t StrideC
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:496
index_t KPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:499
index_t N
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:492
index_t M
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:491
index_t MBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:502
index_t NBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:503
__host__ void Print() const
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:482
index_t BK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:501
index_t MPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:497
index_t K
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:493
index_t StrideA
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:494
index_t AK0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:500
__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:460
index_t NPadded
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:498
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:126
static constexpr auto I6
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:133
static __device__ constexpr bool IsValidCompilationParameter()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:594
static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:565
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:149
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:144
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:713
static constexpr __host__ index_t CalculateKBlockLoopTailNum(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:720
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:532
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:731
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:749
static constexpr auto AK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:139
static constexpr auto I1
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:128
static constexpr auto I2
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:129
static constexpr auto AK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:137
static constexpr auto I0
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:127
static __host__ auto CalculateMBlock(index_t M)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:198
static constexpr auto BK0Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:138
static __host__ auto CalculateAK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:164
static __host__ auto CalculateNBlock(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:203
static constexpr __device__ auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:534
static constexpr auto I4
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:131
static constexpr __device__ auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:542
__host__ static constexpr __device__ auto MakeGemmMmaTileDescriptor(const TileDesc_K0_MN_K1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:209
static constexpr auto I5
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:132
static constexpr __device__ auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:550
static __host__ auto CalculateBK0(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:181
static constexpr auto BK1Number
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:140
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:223
__host__ static constexpr __device__ auto MakeAMmaTileDescriptor_M0_M1_M2_K(const ABlockDesc_AK0_M_AK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:389
static constexpr auto I3
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:130
static __host__ auto CalculateKPadded(index_t K)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:159
__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:406
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:154
__host__ static constexpr __device__ auto MakeBMmaTileDescriptor_N0_N1_N2_K(const BBlockDesc_BK0_N_BK1 &)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:398
static constexpr auto I7
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:134
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:609
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:305
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdl_cshuffle_v2.hpp:142
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