20 template <
typename GridwiseGemm,
23 typename AGridDesc_K0_M_K1,
24 typename BGridDesc_K0_N_K1,
25 typename CGridDesc_M_N,
26 bool HasMainKBlockLoop>
28 #if CK_USE_LAUNCH_BOUNDS
31 #if CK_USE_WAVES_PER_EU
32 __attribute__((amdgpu_waves_per_eu(CK_MIN_WAVES_PER_EU, CK_MAX_WAVES_PER_EU)))
35 const FloatAB* __restrict__ p_b_grid,
36 FloatC* __restrict__ p_c_grid,
37 const AGridDesc_K0_M_K1 a_grid_desc_k0_m_k1,
38 const BGridDesc_K0_N_K1 b_grid_desc_k0_n_k1,
39 const CGridDesc_M_N c_grid_desc_m_n)
41 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
43 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
45 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
47 GridwiseGemm::template Run<HasMainKBlockLoop>(p_a_grid,
59 ignore = a_grid_desc_k0_m_k1;
60 ignore = b_grid_desc_k0_n_k1;
65 template <
typename Gr
idwiseGemm,
bool HasMainKBlockLoop>
67 #if CK_USE_LAUNCH_BOUNDS
70 #if CK_USE_WAVES_PER_EU
71 __attribute__((amdgpu_waves_per_eu(CK_MIN_WAVES_PER_EU, CK_MAX_WAVES_PER_EU)))
75 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
77 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
79 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
81 const auto a_grid_desc_k0_m_k1 =
83 karg.M, karg.MPadded, karg.K, karg.K0, karg.StrideA));
84 const auto b_grid_desc_k0_n_k1 =
86 karg.K, karg.N, karg.NPadded, karg.K0, karg.StrideB));
88 karg.M, karg.MPadded, karg.N, karg.NPadded, karg.StrideC));
90 GridwiseGemm::template Run<HasMainKBlockLoop>(karg.p_a_grid,
108 typename AElementwiseOperation,
109 typename BElementwiseOperation,
110 typename CElementwiseOperation,
119 typename ABlockTransferThreadClusterLengths_K0_M_K1,
120 typename ABlockTransferThreadClusterArrangeOrder,
121 typename ABlockTransferSrcAccessOrder,
122 index_t ABlockTransferSrcVectorDim,
123 index_t ABlockTransferSrcScalarPerVector,
124 index_t ABlockTransferDstScalarPerVector_K1,
125 bool AThreadTransferSrcResetCoordinateAfterRun,
126 bool ABlockLdsExtraM,
127 typename BBlockTransferThreadClusterLengths_K0_N_K1,
128 typename BBlockTransferThreadClusterArrangeOrder,
129 typename BBlockTransferSrcAccessOrder,
130 index_t BBlockTransferSrcVectorDim,
131 index_t BBlockTransferSrcScalarPerVector,
132 index_t BBlockTransferDstScalarPerVector_K1,
133 bool BThreadTransferSrcResetCoordinateAfterRun,
134 bool BBlockLdsExtraN,
135 typename CThreadTransferSrcDstAccessOrder,
136 index_t CThreadTransferSrcDstVectorDim,
137 index_t CThreadTransferDstScalarPerVector,
138 index_t NumGemmKPrefetchStage = 1,
163 selected_mfma.k_per_blk)>{};
174 template <
typename CGr
idDesc_M_N>
177 return std::make_tuple(Block2CTileMap::CalculateGridSize(c_grid_desc_m_n), 1, 1);
221 std::cout <<
"problem {" <<
"M:" <<
M <<
", " <<
"N:" <<
N <<
", " <<
"K:" <<
K <<
", "
223 <<
", " <<
"MP:" <<
MPadded <<
", " <<
"NP:" <<
NPadded <<
", " <<
"K0:" <<
K0
250 :
Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_},
263 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>;
269 #if CK_GFX90A_DENORM_WORKAROUND
277 constexpr
auto max_lds_align =
K1;
280 constexpr
auto a_block_desc_k0_m_k1 = [&]() {
281 if constexpr(ABlockLdsExtraM)
294 return a_block_desc_k0_m_k1;
299 constexpr
auto max_lds_align =
K1;
302 constexpr
auto b_block_desc_k0_n_k1 = [&]() {
303 if constexpr(BBlockLdsExtraN)
316 return b_block_desc_k0_n_k1;
326 constexpr
auto max_lds_align =
K1;
328 constexpr
auto a_block_space_size_aligned =
331 constexpr
auto b_block_space_size_aligned =
334 return (a_block_space_size_aligned + b_block_space_size_aligned) *
342 return ck::tensor_operation::device::IsValidGemmCompilationParameter<
351 CGlobalMemoryDataOperation>();
354 template <
typename AGr
idDesc_K0_M_K1,
typename BGr
idDesc_K0_N_K1,
typename CGr
idDesc_M_N>
355 __host__ __device__
static constexpr
bool
357 const BGridDesc_K0_N_K1& b_grid_desc_k0_n_k1,
358 const CGridDesc_M_N& c_grid_desc_m_n)
361 "wrong! K1 need to be known at compile-time");
363 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
364 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
365 "Invalid tuning param!");
367 const auto M = a_grid_desc_k0_m_k1.GetLength(
I1);
368 const auto N = b_grid_desc_k0_n_k1.GetLength(
I1);
369 const auto K0 = a_grid_desc_k0_m_k1.GetLength(
I0);
371 if(!(M == c_grid_desc_m_n.GetLength(
I0) && N == c_grid_desc_m_n.GetLength(
I1) &&
372 K0 == b_grid_desc_k0_n_k1.GetLength(
I0) &&
K1 == a_grid_desc_k0_m_k1.GetLength(
I2) &&
373 K1 == b_grid_desc_k0_n_k1.GetLength(
I2)))
376 if(!(M % MPerBlock == 0 && N % NPerBlock == 0 && K0 % K0PerBlock == 0))
380 const auto num_k_loop = K0 / K0PerBlock;
382 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
394 "wrong! K1 need to be known at compile-time");
396 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
397 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
398 "Invalid tuning param!");
402 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
415 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
418 template <
typename CGr
idDesc>
419 __host__ __device__
static constexpr
auto
422 constexpr
auto max_lds_align =
K1;
425 constexpr
auto a_block_desc_k0_m_k1 = [&]() {
426 if constexpr(ABlockLdsExtraM)
440 constexpr
auto b_block_desc_k0_n_k1 = [&]() {
441 if constexpr(BBlockLdsExtraN)
454 using BlockwiseGemm =
459 decltype(a_block_desc_k0_m_k1),
460 decltype(b_block_desc_k0_n_k1),
469 return BlockwiseGemm::MakeCGridDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(c_grid_desc_m_n);
475 template <
bool HasMainKBlockLoop,
476 typename AGridDesc_K0_M_K1,
477 typename BGridDesc_K0_N_K1,
478 typename CGridDesc_M_N>
482 void* __restrict__ p_shared,
483 const AGridDesc_K0_M_K1& a_grid_desc_k0_m_k1,
484 const BGridDesc_K0_N_K1& b_grid_desc_k0_n_k1,
485 const CGridDesc_M_N& c_grid_desc_m_n)
487 const auto c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
490 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
491 p_a_grid, a_grid_desc_k0_m_k1.GetElementSpaceSize());
492 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
493 p_b_grid, b_grid_desc_k0_n_k1.GetElementSpaceSize());
494 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
495 p_c_grid, c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetElementSpaceSize());
497 const AElementwiseOperation a_element_op{};
498 const BElementwiseOperation b_element_op{};
499 const CElementwiseOperation c_element_op{};
501 const auto block_2_ctile_map =
505 const auto block_work_idx =
508 if(!block_2_ctile_map.ValidCTileIndex(
510 make_tuple(c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I0),
511 c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I1))))
517 const index_t m_block_data_idx_on_grid =
518 __builtin_amdgcn_readfirstlane(block_work_idx[
I0] * MPerBlock);
520 const index_t n_block_data_idx_on_grid =
521 __builtin_amdgcn_readfirstlane(block_work_idx[
I1] * NPerBlock);
524 constexpr
auto max_lds_align =
K1;
533 auto a_blockwise_copy =
535 AElementwiseOperation,
539 ABlockTransferThreadClusterLengths_K0_M_K1,
540 ABlockTransferThreadClusterArrangeOrder,
543 decltype(a_grid_desc_k0_m_k1),
544 decltype(a_block_desc_k0_m_k1),
545 ABlockTransferSrcAccessOrder,
547 ABlockTransferSrcVectorDim,
549 ABlockTransferSrcScalarPerVector,
550 ABlockTransferDstScalarPerVector_K1,
553 AThreadTransferSrcResetCoordinateAfterRun,
555 NumGemmKPrefetchStage>(
559 a_block_desc_k0_m_k1,
564 auto b_blockwise_copy =
566 BElementwiseOperation,
570 BBlockTransferThreadClusterLengths_K0_N_K1,
571 BBlockTransferThreadClusterArrangeOrder,
574 decltype(b_grid_desc_k0_n_k1),
575 decltype(b_block_desc_k0_n_k1),
576 BBlockTransferSrcAccessOrder,
578 BBlockTransferSrcVectorDim,
580 BBlockTransferSrcScalarPerVector,
581 BBlockTransferDstScalarPerVector_K1,
584 BThreadTransferSrcResetCoordinateAfterRun,
586 NumGemmKPrefetchStage>(
590 b_block_desc_k0_n_k1,
606 decltype(a_block_desc_k0_m_k1),
607 decltype(b_block_desc_k0_n_k1),
617 auto c_thread_buf = blockwise_gemm.GetCThreadBuffer();
620 constexpr
auto a_block_space_size_aligned =
623 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
624 static_cast<ElementDataTypeAB*
>(p_shared), a_block_desc_k0_m_k1.GetElementSpaceSize());
626 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
628 b_block_desc_k0_n_k1.GetElementSpaceSize());
630 constexpr
auto a_block_slice_copy_step =
make_multi_index(K0PerBlock, 0, 0);
631 constexpr
auto b_block_slice_copy_step =
make_multi_index(K0PerBlock, 0, 0);
634 const auto K0 = a_grid_desc_k0_m_k1.GetLength(
I0);
635 const index_t num_k_block_main_loop = __builtin_amdgcn_readfirstlane(K0 / K0PerBlock);
637 GridwiseGemmPipe::template Run<HasMainKBlockLoop>(a_grid_desc_k0_m_k1,
638 a_block_desc_k0_m_k1,
642 a_block_slice_copy_step,
644 b_block_desc_k0_n_k1,
648 b_block_slice_copy_step,
651 num_k_block_main_loop);
655 constexpr
auto c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
656 blockwise_gemm.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
658 constexpr
auto c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2 =
659 blockwise_gemm.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
661 constexpr
auto M0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I0);
662 constexpr
auto N0 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I1);
663 constexpr
auto M1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I2);
664 constexpr
auto N1 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I3);
665 constexpr
auto M2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I4);
666 constexpr
auto M3 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I5);
667 constexpr
auto M4 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I6);
668 constexpr
auto N2 = c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2.GetLength(
I7);
672 const auto c_thread_mtx_on_block =
673 blockwise_gemm.CalculateCThreadOriginDataIndex(
I0,
I0,
I0,
I0);
675 const index_t m_thread_data_on_grid =
676 m_block_data_idx_on_grid + c_thread_mtx_on_block[
I0];
678 const index_t n_thread_data_on_grid =
679 n_block_data_idx_on_grid + c_thread_mtx_on_block[
I1];
681 const auto m_thread_data_on_grid_to_m0_m1_m2_m3_m4_adaptor =
687 const auto m_thread_data_on_grid_idx =
688 m_thread_data_on_grid_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
696 const auto n_thread_data_on_grid_idx =
697 n_thread_data_on_grid_to_n0_n1_n2_adaptor.CalculateBottomIndex(
703 decltype(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2),
704 decltype(c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2),
705 CElementwiseOperation,
707 CThreadTransferSrcDstAccessOrder,
708 CThreadTransferSrcDstVectorDim,
709 CThreadTransferDstScalarPerVector,
710 CGlobalMemoryDataOperation,
713 c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2,
715 n_thread_data_on_grid_idx[
I0],
716 m_thread_data_on_grid_idx[
I1],
717 n_thread_data_on_grid_idx[
I1],
718 m_thread_data_on_grid_idx[
I2],
719 m_thread_data_on_grid_idx[
I3],
720 m_thread_data_on_grid_idx[
I4],
721 n_thread_data_on_grid_idx[
I2]),
724 c_thread_copy.Run(c_thread_desc_m0_n0_m1_n1_m2_m3_m4_n2,
727 c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2,
741 typename AElementwiseOperation,
742 typename BElementwiseOperation,
743 typename CElementwiseOperation,
753 typename ABlockTransferThreadClusterLengths_K0_M_K1,
754 typename ABlockTransferThreadClusterArrangeOrder,
755 typename ABlockTransferSrcAccessOrder,
756 index_t ABlockTransferSrcVectorDim,
757 index_t ABlockTransferSrcScalarPerVector,
758 index_t ABlockTransferDstScalarPerVector_K1,
759 bool AThreadTransferSrcResetCoordinateAfterRun,
760 bool ABlockLdsExtraM,
761 typename BBlockTransferThreadClusterLengths_K0_N_K1,
762 typename BBlockTransferThreadClusterArrangeOrder,
763 typename BBlockTransferSrcAccessOrder,
764 index_t BBlockTransferSrcVectorDim,
765 index_t BBlockTransferSrcScalarPerVector,
766 index_t BBlockTransferDstScalarPerVector_K1,
767 bool BThreadTransferSrcResetCoordinateAfterRun,
768 bool BBlockLdsExtraN,
769 typename CThreadTransferSrcDstAccessOrder,
770 index_t CThreadTransferSrcDstVectorDim,
771 index_t CThreadTransferDstScalarPerVector,
772 index_t NumGemmKPrefetchStage = 1,
780 CGlobalMemoryDataOperation,
781 AElementwiseOperation,
782 BElementwiseOperation,
783 CElementwiseOperation,
792 ABlockTransferThreadClusterLengths_K0_M_K1,
793 ABlockTransferThreadClusterArrangeOrder,
794 ABlockTransferSrcAccessOrder,
795 ABlockTransferSrcVectorDim,
796 ABlockTransferSrcScalarPerVector,
797 ABlockTransferDstScalarPerVector_K1,
798 AThreadTransferSrcResetCoordinateAfterRun,
800 BBlockTransferThreadClusterLengths_K0_N_K1,
801 BBlockTransferThreadClusterArrangeOrder,
802 BBlockTransferSrcAccessOrder,
803 BBlockTransferSrcVectorDim,
804 BBlockTransferSrcScalarPerVector,
805 BBlockTransferDstScalarPerVector_K1,
806 BThreadTransferSrcResetCoordinateAfterRun,
808 CThreadTransferSrcDstAccessOrder,
809 CThreadTransferSrcDstVectorDim,
810 CThreadTransferDstScalarPerVector,
811 NumGemmKPrefetchStage,
820 CGlobalMemoryDataOperation,
821 AElementwiseOperation,
822 BElementwiseOperation,
823 CElementwiseOperation,
832 ABlockTransferThreadClusterLengths_K0_M_K1,
833 ABlockTransferThreadClusterArrangeOrder,
834 ABlockTransferSrcAccessOrder,
835 ABlockTransferSrcVectorDim,
836 ABlockTransferSrcScalarPerVector,
837 ABlockTransferDstScalarPerVector_K1,
838 AThreadTransferSrcResetCoordinateAfterRun,
840 BBlockTransferThreadClusterLengths_K0_N_K1,
841 BBlockTransferThreadClusterArrangeOrder,
842 BBlockTransferSrcAccessOrder,
843 BBlockTransferSrcVectorDim,
844 BBlockTransferSrcScalarPerVector,
845 BBlockTransferDstScalarPerVector_K1,
846 BThreadTransferSrcResetCoordinateAfterRun,
848 CThreadTransferSrcDstAccessOrder,
849 CThreadTransferSrcDstVectorDim,
850 CThreadTransferDstScalarPerVector,
851 NumGemmKPrefetchStage,
862 __device__
static auto
865 const auto a_grid_desc_m_k = [&]() {
879 const auto KPad = K0Pad * K1Value;
914 __device__
static auto
917 const auto b_grid_desc_k_n = [&]() {
931 const auto KPad = K0Pad * K1Value;
967 __device__
static auto
970 const auto c_grid_desc_m_n = [&]() {
1006 "wrong! K1 need to be known at compile-time");
1008 static_assert((MPerBlock % (MPerXdl * MXdlPerWave) == 0) &&
1009 (NPerBlock % (NXdlPerWave * NPerXdl)) == 0,
1010 "Invalid tuning param!");
1017 if(!(problem.M % MPerBlock == 0))
1028 if(!(problem.N % NPerBlock == 0))
1039 if(!(problem.K0 % K0PerBlock == 0))
1047 if(problem.K % ABlockTransferSrcScalarPerVector != 0)
1054 if(problem.M % ABlockTransferSrcScalarPerVector != 0)
1062 if(problem.N % BBlockTransferSrcScalarPerVector != 0)
1069 if(problem.K % BBlockTransferSrcScalarPerVector != 0)
1078 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:31
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:30
#define IS_VALID_COMPILATION_PARAMETER_IMPL(CDataType_)
Definition: device_base.hpp:178
__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__ 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
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__host__ constexpr __device__ auto make_naive_tensor_descriptor_aligned(const Tuple< Lengths... > &lengths, Align align)
Definition: tensor_descriptor_helper.hpp:132
__host__ constexpr __device__ auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:425
ushort bhalf_t
Definition: data_type.hpp:30
constexpr detail::ignore_t ignore
Definition: ignore.hpp:20
__device__ uint32_t amd_wave_read_first_lane(uint32_t value)
Definition: amd_wave_read_first_lane.hpp:100
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:47
typename conditional< predicate, X, Y >::type conditional_t
Definition: functional.hpp:115
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__global__ void kernel_gemm_xdlops_v2r3(const FloatAB *__restrict__ p_a_grid, const FloatAB *__restrict__ p_b_grid, FloatC *__restrict__ p_c_grid, const AGridDesc_K0_M_K1 a_grid_desc_k0_m_k1, const BGridDesc_K0_N_K1 b_grid_desc_k0_n_k1, const CGridDesc_M_N c_grid_desc_m_n)
Definition: gridwise_gemm_xdlops_v2r3.hpp:34
__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
PipelineVersion
Definition: gridwise_gemm_pipeline_selector.hpp:18
typename remove_cv< T >::type remove_cv_t
Definition: type.hpp:295
constexpr LoopScheduler make_default_loop_scheduler()
Definition: loop_scheduler.hpp:20
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
Definition: blockwise_gemm_smfmac_xdlops.hpp:44
Definition: gridwise_gemm_xdlops_v2r3.hpp:240
__host__ Argument(const ElementDataTypeAB *p_a_grid_, const ElementDataTypeAB *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_xdlops_v2r3.hpp:241
const ElementDataTypeAB * p_b_grid
Definition: gridwise_gemm_xdlops_v2r3.hpp:258
const ElementDataTypeAB * p_a_grid
Definition: gridwise_gemm_xdlops_v2r3.hpp:257
FloatC * p_c_grid
Definition: gridwise_gemm_xdlops_v2r3.hpp:259
Definition: gridwise_gemm_xdlops_v2r3.hpp:200
index_t NPadded
Definition: gridwise_gemm_xdlops_v2r3.hpp:234
index_t K
Definition: gridwise_gemm_xdlops_v2r3.hpp:229
__host__ Problem(index_t M_, index_t N_, index_t K_, index_t StrideA_, index_t StrideB_, index_t StrideC_)
Definition: gridwise_gemm_xdlops_v2r3.hpp:201
index_t StrideB
Definition: gridwise_gemm_xdlops_v2r3.hpp:231
index_t N
Definition: gridwise_gemm_xdlops_v2r3.hpp:228
index_t StrideC
Definition: gridwise_gemm_xdlops_v2r3.hpp:232
__host__ void Print() const
Definition: gridwise_gemm_xdlops_v2r3.hpp:219
index_t StrideA
Definition: gridwise_gemm_xdlops_v2r3.hpp:230
index_t MPadded
Definition: gridwise_gemm_xdlops_v2r3.hpp:233
index_t M
Definition: gridwise_gemm_xdlops_v2r3.hpp:227
index_t K0
Definition: gridwise_gemm_xdlops_v2r3.hpp:235
Definition: gridwise_gemm_xdlops_v2r3.hpp:814
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdlops_v2r3.hpp:1003
static __device__ auto MakeCGridDescriptor_M_N(index_t M, index_t MPad, index_t N, index_t NPad, index_t StrideC)
Definition: gridwise_gemm_xdlops_v2r3.hpp:968
static constexpr auto I1
Definition: gridwise_gemm_xdlops_v2r3.hpp:144
static __device__ auto MakeBGridDescriptor_K0_N_K1(index_t K, index_t N, index_t NPad, index_t K0, index_t StrideB)
Definition: gridwise_gemm_xdlops_v2r3.hpp:915
static __device__ auto MakeAGridDescriptor_K0_M_K1(index_t M, index_t MPad, index_t K, index_t K0, index_t StrideA)
Definition: gridwise_gemm_xdlops_v2r3.hpp:863
static constexpr auto K1
Definition: gridwise_gemm_xdlops_v2r3.hpp:160
Definition: gridwise_gemm_xdlops_v2r3.hpp:142
static constexpr __host__ bool CheckValidity(const Problem &problem)
Definition: gridwise_gemm_xdlops_v2r3.hpp:391
static __host__ auto CalculateK0(index_t K)
Definition: gridwise_gemm_xdlops_v2r3.hpp:196
static constexpr auto I5
Definition: gridwise_gemm_xdlops_v2r3.hpp:148
__host__ static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdlops_v2r3.hpp:319
static constexpr auto I2
Definition: gridwise_gemm_xdlops_v2r3.hpp:145
static __device__ constexpr bool IsValidCompilationParameter()
Definition: gridwise_gemm_xdlops_v2r3.hpp:340
static __device__ void Run(const ElementDataTypeAB *p_a_grid, const ElementDataTypeAB *p_b_grid, FloatC *p_c_grid, void *__restrict__ p_shared, const AGridDesc_K0_M_K1 &a_grid_desc_k0_m_k1, const BGridDesc_K0_N_K1 &b_grid_desc_k0_n_k1, const CGridDesc_M_N &c_grid_desc_m_n)
Definition: gridwise_gemm_xdlops_v2r3.hpp:479
static constexpr auto is_scale_mfma
Definition: gridwise_gemm_xdlops_v2r3.hpp:159
static constexpr auto I6
Definition: gridwise_gemm_xdlops_v2r3.hpp:149
conditional_t< is_same_v< FloatAB, ck::tf32_t >, float, FloatAB > ElementDataTypeAB
Definition: gridwise_gemm_xdlops_v2r3.hpp:167
static constexpr auto I1
Definition: gridwise_gemm_xdlops_v2r3.hpp:144
static __host__ auto CalculateGridSize(const CGridDesc_M_N &c_grid_desc_m_n)
Definition: gridwise_gemm_xdlops_v2r3.hpp:175
__host__ static constexpr __device__ auto GetBBlockDescriptor_K0PerBlock_NPerBlock_K1()
Definition: gridwise_gemm_xdlops_v2r3.hpp:297
static __host__ auto CalculateMPadded(index_t M)
Definition: gridwise_gemm_xdlops_v2r3.hpp:186
static constexpr __host__ bool CalculateHasMainKBlockLoop(index_t K)
Definition: gridwise_gemm_xdlops_v2r3.hpp:411
static constexpr auto I7
Definition: gridwise_gemm_xdlops_v2r3.hpp:150
FloatAB FloatABAdjusted
Definition: gridwise_gemm_xdlops_v2r3.hpp:272
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdlops_v2r3.hpp:169
static constexpr auto I3
Definition: gridwise_gemm_xdlops_v2r3.hpp:146
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage, LoopSched >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdlops_v2r3.hpp:263
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdlops_v2r3.hpp:165
static __host__ auto CalculateNPadded(index_t N)
Definition: gridwise_gemm_xdlops_v2r3.hpp:191
__host__ static constexpr __device__ auto MakeCGridDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(const CGridDesc &c_grid_desc_m_n)
Definition: gridwise_gemm_xdlops_v2r3.hpp:420
__host__ static constexpr __device__ auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1()
Definition: gridwise_gemm_xdlops_v2r3.hpp:275
static constexpr auto I4
Definition: gridwise_gemm_xdlops_v2r3.hpp:147
static __host__ auto CalculateGridSize(index_t M, index_t N)
Definition: gridwise_gemm_xdlops_v2r3.hpp:181
static constexpr auto I0
Definition: gridwise_gemm_xdlops_v2r3.hpp:143
__host__ static constexpr __device__ bool CheckValidity(const AGridDesc_K0_M_K1 &a_grid_desc_k0_m_k1, const BGridDesc_K0_N_K1 &b_grid_desc_k0_n_k1, const CGridDesc_M_N &c_grid_desc_m_n)
Definition: gridwise_gemm_xdlops_v2r3.hpp:356
static constexpr bool is_single_rate_mfma
Definition: gridwise_gemm_xdlops_v2r3.hpp:153
static constexpr auto K1
Definition: gridwise_gemm_xdlops_v2r3.hpp:160
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
Blockwise data transfer.
Definition: thread_group_tensor_slice_transfer_v4r1.hpp:46
Definition: threadwise_tensor_slice_transfer.hpp:39
Definition: integral_constant.hpp:20
Definition: is_known_at_compile_time.hpp:14
Definition: device_base.hpp:197
Definition: unary_element_wise_operation.hpp:334