23 template <
typename LowLengths>
49 static_assert(LowerIndex::Size() == NDimLow,
"wrong!");
56 __host__ __device__ constexpr
const auto&
GetUpperLengths()
const {
return up_lengths_; }
58 template <
typename LowIdx,
typename UpIdx>
60 const UpIdx& idx_up)
const
62 static_assert(LowIdx::Size() == NDimLow && UpIdx::Size() == 1,
63 "wrong! inconsistent # of dimension");
69 idx_low(i) = tmp / this->low_lengths_scan_[i];
70 tmp %= this->low_lengths_scan_[i];
76 template <
typename LowIdxDiff,
82 const UpIdxDiff& idx_up_diff,
84 const UpIdx& idx_up_new,
87 static_assert(LowIdxDiff::Size() == NDimLow && UpIdxDiff::Size() == 1 &&
88 LowIdx::Size() == NDimLow && UpIdx::Size() == 1,
89 "wrong! inconsistent # of dimension");
92 constexpr
auto INm1 =
Number<NDimLow - 1>{};
97 idx_diff_low(INm1) = idx_up_diff[I0];
114 template <
typename UpIdx>
115 __host__ __device__
static constexpr
bool
121 __host__ __device__
void Print()
const
124 printf(
"Merge_v3_direct_division_mod_wrw, ");
125 printf(
"low_lengths_ ");
127 printf(
"low_lengths_scan_ ");
129 printf(
"up_lengths_ ");
135 template <
typename LowLengths>
141 template <
typename GridwiseGemm,
145 typename AGridDesc_B_K0_M_K1,
146 typename BGridDesc_B_K0_N_K1,
147 typename CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
148 typename AElementwiseOperation,
149 typename BElementwiseOperation,
150 typename CElementwiseOperation,
151 typename CBlockClusterAdaptor,
152 bool HasMainKBlockLoop>
154 #if CK_USE_LAUNCH_BOUNDS
158 const FloatB* __restrict__ p_b_grid,
159 FloatC* __restrict__ p_c_grid,
160 const AGridDesc_B_K0_M_K1 a_b_k0_m_k1_grid_desc,
161 const BGridDesc_B_K0_N_K1 b_b_k0_n_k1_grid_desc,
162 const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock
163 c_grid_desc_mblock_mperblock_nblock_nperblock,
164 const AElementwiseOperation a_element_op,
165 const BElementwiseOperation b_element_op,
166 const CElementwiseOperation c_element_op,
167 const CBlockClusterAdaptor c_block_cluster_adaptor)
169 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx11__) || \
171 if constexpr(GridwiseGemm::template IsValidCompilationParameter<>())
173 __shared__
char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
175 GridwiseGemm::template Run<HasMainKBlockLoop>(p_a_grid,
179 a_b_k0_m_k1_grid_desc,
180 b_b_k0_n_k1_grid_desc,
181 c_grid_desc_mblock_mperblock_nblock_nperblock,
185 c_block_cluster_adaptor);
191 ignore = a_b_k0_m_k1_grid_desc;
192 ignore = b_b_k0_n_k1_grid_desc;
193 ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
197 ignore = c_block_cluster_adaptor;
207 typename AGridDesc_B_K0_M_K1,
208 typename BGridDesc_B_K0_N_K1,
209 typename CMNGridDesc,
210 typename AElementwiseOperation,
211 typename BElementwiseOperation,
212 typename CElementwiseOperation,
221 typename ABlockTransferThreadClusterLengths_K0_M_K1,
222 typename ABlockTransferThreadClusterArrangeOrder,
223 typename ABlockTransferSrcAccessOrder,
224 index_t ABlockTransferSrcVectorDim,
225 index_t ABlockTransferSrcScalarPerVector,
226 index_t ABlockTransferDstScalarPerVector_K1,
227 bool AThreadTransferSrcResetCoordinateAfterRun,
228 bool ABlockLdsExtraM,
232 typename BBlockTransferThreadClusterLengths_K0_N_K1,
233 typename BBlockTransferThreadClusterArrangeOrder,
234 typename BBlockTransferSrcAccessOrder,
235 index_t BBlockTransferSrcVectorDim,
236 index_t BBlockTransferSrcScalarPerVector,
237 index_t BBlockTransferDstScalarPerVector_K1,
238 bool BThreadTransferSrcResetCoordinateAfterRun,
239 bool BBlockLdsExtraN,
243 index_t CShuffleMRepeatPerShuffle,
244 index_t CShuffleNRepeatPerShuffle,
245 index_t CBlockTransferScalarPerVector_NWaveNPerXDL,
246 typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
247 bool ABlockLdsExtraM1Wrw =
false,
248 bool BBlockLdsExtraN1Wrw =
false,
249 index_t NumGemmKPrefetchStage = 1,
251 typename ComputeTypeA = FloatA,
252 typename ComputeTypeB = ComputeTypeA>
270 decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage>())>;
277 #if CK_GFX90A_DENORM_WORKAROUND
299 constexpr
auto max_lds_align = K1;
302 constexpr
auto a_block_desc_k0_m_k1 = [&]() {
303 if constexpr(ABlockLdsExtraM)
305 if constexpr(ABlockLdsExtraM1Wrw)
316 a_block_desc_k0_m0_m1_k1,
324 return a_block_desc_k0_m_k1_tmp;
340 return a_block_desc_k0_m_k1;
345 constexpr
auto max_lds_align = K1;
348 constexpr
auto a_block_desc_b_k0_m_k1 = [&]() {
349 if constexpr(ABlockLdsExtraM)
351 if constexpr(ABlockLdsExtraM1Wrw)
367 a_block_desc_b_k0_m0_m1_k1,
376 return a_block_desc_b_k0_m_k1_tmp;
396 return a_block_desc_b_k0_m_k1;
401 constexpr
auto max_lds_align = K1;
404 constexpr
auto b_block_desc_k0_n_k1 = [&]() {
405 if constexpr(BBlockLdsExtraN)
407 if constexpr(BBlockLdsExtraN1Wrw)
418 b_block_desc_k0_n0_n1_k1,
426 return b_block_desc_k0_n_k1_tmp;
442 return b_block_desc_k0_n_k1;
447 constexpr
auto max_lds_align = K1;
450 constexpr
auto b_block_desc_b_k0_n_k1 = [&]() {
451 if constexpr(BBlockLdsExtraN)
453 if constexpr(BBlockLdsExtraN1Wrw)
469 b_block_desc_b_k0_n0_n1_k1,
478 return b_block_desc_b_k0_n_k1_tmp;
498 return b_block_desc_b_k0_n_k1;
503 constexpr
auto max_lds_align = K1;
506 constexpr
auto a_b_k0_m_k1_block_desc = GetABlockDescriptor_Batch_K0PerBlock_MPerBlock_K1();
509 constexpr
auto b_b_k0_n_k1_block_desc = GetBBlockDescriptor_Batch_K0PerBlock_NPerBlock_K1();
513 a_b_k0_m_k1_block_desc.GetElementSpaceSize(), max_lds_align);
516 b_b_k0_n_k1_block_desc.GetElementSpaceSize(), max_lds_align);
518 constexpr
auto c_block_size =
519 GetCBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock().GetElementSpaceSize();
523 c_block_size *
sizeof(FloatC));
530 return ck::tensor_operation::device::IsValidGemmCompilationParameter<
539 CGlobalMemoryDataOperation>();
542 template <
typename Block2CTileMap>
543 __host__ __device__
static constexpr
bool
545 const BGridDesc_B_K0_N_K1& b_b_k0_n_k1_grid_desc,
546 const CMNGridDesc& c_m_n_grid_desc,
547 const Block2CTileMap& block_2_ctile_map)
550 "wrong! K1 need to be known at compile-time");
552 static_assert((MPerBlock % (MPerXdl * MRepeat) == 0) &&
553 (NPerBlock % (NRepeat * NPerXdl)) == 0,
554 "Invalid tuning param!");
556 const auto M = a_b_k0_m_k1_grid_desc.GetLength(I2);
557 const auto N = b_b_k0_n_k1_grid_desc.GetLength(I2);
558 const auto K0 = a_b_k0_m_k1_grid_desc.GetLength(I1);
559 const auto KBatch = a_b_k0_m_k1_grid_desc.GetLength(I0);
562 const auto num_k_loop = K0 / K0PerBlock;
564 if(!GridwiseGemmPipe::IsSupported(num_k_loop))
569 if(!(M == c_m_n_grid_desc.GetLength(I0) && N == c_m_n_grid_desc.GetLength(I1) &&
570 K0 == b_b_k0_n_k1_grid_desc.GetLength(I1) &&
571 K1 == a_b_k0_m_k1_grid_desc.GetLength(I3) &&
572 K1 == b_b_k0_n_k1_grid_desc.GetLength(I3) &&
573 KBatch == b_b_k0_n_k1_grid_desc.GetLength(I0)))
576 if(!(M % MPerBlock == 0 && N % NPerBlock == 0 && K0 % K0PerBlock == 0))
579 if(!block_2_ctile_map.CheckValidity(c_m_n_grid_desc))
591 const index_t num_loop = K0 / K0PerBlock;
593 return GridwiseGemmPipe::CalculateHasMainLoop(num_loop);
598 __host__ __device__
static constexpr
auto
601 const auto M = c_m_n_grid_desc.GetLength(I0);
602 const auto N = c_m_n_grid_desc.GetLength(I1);
604 const auto MBlock = M / MPerBlock;
605 const auto NBlock = N / NPerBlock;
620 c_m_n_grid_desc, M01, N01, KBatch);
623 __host__ __device__
static constexpr
auto
626 constexpr
index_t MWave = MPerBlock / (MRepeat * MPerXdl);
627 constexpr
index_t NWave = NPerBlock / (NRepeat * NPerXdl);
637 decltype(MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(CMNGridDesc{}));
640 template <
bool HasMainKBlockLoop>
641 __device__
static void Run(
const FloatA* __restrict__ p_a_grid,
642 const FloatB* __restrict__ p_b_grid,
643 FloatC* __restrict__ p_c_grid,
644 void* __restrict__ p_shared,
645 const AGridDesc_B_K0_M_K1& a_b_k0_m_k1_grid_desc,
646 const BGridDesc_B_K0_N_K1& b_b_k0_n_k1_grid_desc,
648 c_grid_desc_mblock_mperblock_nblock_nperblock,
649 const AElementwiseOperation& a_element_op,
650 const BElementwiseOperation& b_element_op,
651 const CElementwiseOperation& c_element_op,
654 const auto a_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
655 p_a_grid, a_b_k0_m_k1_grid_desc.GetElementSpaceSize());
656 const auto b_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
657 p_b_grid, b_b_k0_n_k1_grid_desc.GetElementSpaceSize());
658 auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
659 p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
661 const auto K0 = a_b_k0_m_k1_grid_desc.GetLength(I1);
664 const auto block_work_idx =
667 const index_t k_batch_id = block_work_idx[I0];
669 if(!c_block_cluster_adaptor.ValidCTileIndex(
670 make_tuple(block_work_idx[I1], block_work_idx[I2]),
671 make_tuple(c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(I0),
672 c_grid_desc_mblock_mperblock_nblock_nperblock.GetLength(I2))))
678 const index_t m_block_data_idx_on_grid =
679 __builtin_amdgcn_readfirstlane(block_work_idx[I1] * MPerBlock);
681 const index_t n_block_data_idx_on_grid =
682 __builtin_amdgcn_readfirstlane(block_work_idx[I2] * NPerBlock);
685 constexpr
auto max_lds_align = K1;
688 constexpr
auto a_k0_m_k1_block_desc = GetABlockDescriptor_K0PerBlock_MPerBlock_K1();
690 constexpr
auto a_b_k0_m_k1_block_desc = GetABlockDescriptor_Batch_K0PerBlock_MPerBlock_K1();
692 constexpr
auto b_k0_n_k1_block_desc = GetBBlockDescriptor_K0PerBlock_NPerBlock_K1();
694 constexpr
auto b_b_k0_n_k1_block_desc = GetBBlockDescriptor_Batch_K0PerBlock_NPerBlock_K1();
696 auto a_blockwise_copy =
698 AElementwiseOperation,
700 InMemoryDataOperationEnum::Set,
702 ABlockTransferThreadClusterLengths_K0_M_K1,
703 ABlockTransferThreadClusterArrangeOrder,
706 decltype(a_b_k0_m_k1_grid_desc),
707 decltype(a_b_k0_m_k1_block_desc),
708 ABlockTransferSrcAccessOrder,
710 ABlockTransferSrcVectorDim,
712 ABlockTransferSrcScalarPerVector,
713 ABlockTransferDstScalarPerVector_K1,
716 AThreadTransferSrcResetCoordinateAfterRun,
718 a_b_k0_m_k1_grid_desc,
721 a_b_k0_m_k1_block_desc,
726 auto b_blockwise_copy =
728 BElementwiseOperation,
730 InMemoryDataOperationEnum::Set,
732 BBlockTransferThreadClusterLengths_K0_N_K1,
733 BBlockTransferThreadClusterArrangeOrder,
736 decltype(b_b_k0_n_k1_grid_desc),
737 decltype(b_b_k0_n_k1_block_desc),
738 BBlockTransferSrcAccessOrder,
740 BBlockTransferSrcVectorDim,
742 BBlockTransferSrcScalarPerVector,
743 BBlockTransferDstScalarPerVector_K1,
746 BThreadTransferSrcResetCoordinateAfterRun,
748 b_b_k0_n_k1_grid_desc,
751 b_b_k0_n_k1_block_desc,
762 constexpr
bool is_single_rate_mfma =
770 constexpr
auto is_scale_mfma =
false;
777 is_scale_mfma>::selected_mfma.k_per_blk);
779 auto blockwise_gemm =
784 decltype(a_k0_m_k1_block_desc),
785 decltype(b_k0_n_k1_block_desc),
795 constexpr
auto a_block_space_size =
798 constexpr
auto a_block_slice_copy_step =
make_multi_index(0, K0PerBlock, 0, 0);
799 constexpr
auto b_block_slice_copy_step =
make_multi_index(0, K0PerBlock, 0, 0);
801 auto a_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
802 static_cast<FloatAAdjusted*
>(p_shared), a_k0_m_k1_block_desc.GetElementSpaceSize());
804 auto b_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
806 b_k0_n_k1_block_desc.GetElementSpaceSize());
809 const index_t K0BlockMainLoop = __builtin_amdgcn_readfirstlane(K0 / K0PerBlock);
811 GridwiseGemmPipe::template Run<HasMainKBlockLoop>(a_b_k0_m_k1_grid_desc,
812 a_b_k0_m_k1_block_desc,
816 a_block_slice_copy_step,
817 b_b_k0_n_k1_grid_desc,
818 b_b_k0_n_k1_block_desc,
822 b_block_slice_copy_step,
829 constexpr
index_t MWave = MPerBlock / (MRepeat * MPerXdl);
830 constexpr
index_t NWave = NPerBlock / (NRepeat * NPerXdl);
832 constexpr
auto c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc =
833 blockwise_gemm.GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
835 constexpr
auto c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc =
836 blockwise_gemm.GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2();
838 constexpr
auto M0 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I0);
839 constexpr
auto N0 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I1);
840 constexpr
auto M1 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I2);
841 constexpr
auto N1 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I3);
842 constexpr
auto M2 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I4);
843 constexpr
auto M3 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I5);
844 constexpr
auto M4 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I6);
845 constexpr
auto N2 = c_m0_n0_m1_n1_m2_m3_m4_n2_block_desc.GetLength(I7);
847 constexpr
auto c_block_desc_mblock_mperblock_nblock_nperblock =
848 GetCBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock();
850 auto c_block_buf = make_dynamic_buffer<AddressSpaceEnum::Lds>(
851 static_cast<FloatC*
>(p_shared),
852 c_block_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
854 static_assert(M1 == MWave,
"");
855 static_assert(N1 == NWave,
"");
856 static_assert(M2 * M3 * M4 == MPerXdl,
"");
857 static_assert(N2 == NPerXdl,
"");
860 c_block_desc_mblock_mperblock_nblock_nperblock,
878 const auto c_thread_mtx_on_block =
879 blockwise_gemm.CalculateCThreadOriginDataIndex(I0, I0, I0, I0);
881 const index_t m_thread_data_on_block = c_thread_mtx_on_block[I0];
882 const index_t n_thread_data_on_block = c_thread_mtx_on_block[I1];
884 const auto m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor =
890 const auto m_thread_data_on_block_idx =
891 m_thread_data_on_block_to_m0_m1_m2_m3_m4_adaptor.CalculateBottomIndex(
894 const auto n_thread_data_on_block_to_n0_n1_n2_adaptor =
900 const auto n_thread_data_on_block_idx =
901 n_thread_data_on_block_to_n0_n1_n2_adaptor.CalculateBottomIndex(
905 auto c_thread_copy_vgpr_to_lds =
908 decltype(c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc),
909 decltype(c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2),
912 CShuffleNRepeatPerShuffle,
922 InMemoryDataOperationEnum::Set,
925 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
928 m_thread_data_on_block_idx[I1],
929 n_thread_data_on_block_idx[I1],
930 m_thread_data_on_block_idx[I2],
931 m_thread_data_on_block_idx[I3],
932 m_thread_data_on_block_idx[I4],
933 n_thread_data_on_block_idx[I2]),
939 CElementwiseOperation,
940 CGlobalMemoryDataOperation,
942 CShuffleMRepeatPerShuffle * MWave * MPerXdl,
944 CShuffleNRepeatPerShuffle * NWave * NPerXdl>,
945 CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
949 decltype(c_block_desc_mblock_mperblock_nblock_nperblock),
950 decltype(c_grid_desc_mblock_mperblock_nblock_nperblock),
953 CBlockTransferScalarPerVector_NWaveNPerXDL,
956 {c_block_desc_mblock_mperblock_nblock_nperblock,
958 c_grid_desc_mblock_mperblock_nblock_nperblock,
962 constexpr
auto mxdlperwave_forward_step =
964 constexpr
auto nxdlperwave_forward_step =
966 constexpr
auto nxdlperwave_backward_step =
970 constexpr
auto mxdlperwave = mxdlperwave_iter;
973 constexpr
bool nxdlperwave_forward_sweep =
974 (mxdlperwave % (2 * CShuffleMRepeatPerShuffle) == 0);
976 constexpr
index_t nxdlperwave_value =
977 nxdlperwave_forward_sweep
979 : (NRepeat - nxdlperwave_iter - CShuffleNRepeatPerShuffle);
987 c_thread_copy_vgpr_to_lds.Run(
988 c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc,
989 make_tuple(mxdlperwave, nxdlperwave, I0, I0, I0, I0, I0, I0),
991 c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2,
998 c_block_copy_lds_to_global.Run(c_block_desc_mblock_mperblock_nblock_nperblock,
1000 c_grid_desc_mblock_mperblock_nblock_nperblock,
1004 if constexpr(nxdlperwave_forward_sweep &&
1005 (nxdlperwave < NRepeat - CShuffleNRepeatPerShuffle))
1007 c_block_copy_lds_to_global.MoveDstSliceWindow(
1008 c_grid_desc_mblock_mperblock_nblock_nperblock,
1009 nxdlperwave_forward_step);
1011 else if constexpr((!nxdlperwave_forward_sweep) && (nxdlperwave > 0))
1013 c_block_copy_lds_to_global.MoveDstSliceWindow(
1014 c_grid_desc_mblock_mperblock_nblock_nperblock,
1015 nxdlperwave_backward_step);
1020 if constexpr(mxdlperwave < MRepeat - CShuffleMRepeatPerShuffle)
1022 c_block_copy_lds_to_global.MoveDstSliceWindow(
1023 c_grid_desc_mblock_mperblock_nblock_nperblock, mxdlperwave_forward_step);
CK_TILE_DEVICE void block_sync_lds()
Definition: arch.hpp:192
#define CK_MIN_BLOCK_PER_CU
Definition: ck.hpp:31
#define CK_MAX_THREAD_PER_BLOCK
Definition: ck.hpp:30
__host__ constexpr __device__ auto integer_least_multiple(X x, Y y)
Definition: math.hpp:78
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
__host__ __device__ multiplies() -> multiplies< void, void >
FIXME: create macro to replace 'host device' and nothing more.
__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
__global__ void kernel_gemm_xdlops_bwd_weight(const FloatA *__restrict__ p_a_grid, const FloatB *__restrict__ p_b_grid, FloatC *__restrict__ p_c_grid, const AGridDesc_B_K0_M_K1 a_b_k0_m_k1_grid_desc, const BGridDesc_B_K0_N_K1 b_b_k0_n_k1_grid_desc, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock, const AElementwiseOperation a_element_op, const BElementwiseOperation b_element_op, const CElementwiseOperation c_element_op, const CBlockClusterAdaptor c_block_cluster_adaptor)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:157
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_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
__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
typename conditional< predicate, X, Y >::type conditional_t
Definition: functional.hpp:115
__host__ constexpr __device__ auto container_reverse_exclusive_scan(const Array< TData, NSize > &x, Reduce f, TData init)
Definition: container_helper.hpp:213
__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
int32_t index_t
Definition: ck.hpp:299
__host__ constexpr __device__ auto container_reduce(const Container &x, Reduce reduce, Init init, Number< IBegin >=Number< 0 >{}, Number< IEnd >=Number< Container::Size()>{}, Number< IStep >=Number< 1 >{})
Definition: container_helper.hpp:111
__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_descriptor.hpp:319
PipelineVersion
Definition: gridwise_gemm_pipeline_selector.hpp:18
__host__ __device__ void print_multi_index(const Tuple< Xs... > &x)
Definition: statically_indexed_array_multi_index.hpp:147
typename remove_cv< T >::type remove_cv_t
Definition: type.hpp:295
__host__ constexpr __device__ auto make_merge_transform_v4_no_carry(const LowLengths &low_lengths)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:136
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
Definition: block_to_ctile_map.hpp:720
Definition: blockwise_gemm_smfmac_xdlops.hpp:44
__host__ constexpr __device__ auto & GetCThreadBuffer()
Definition: blockwise_gemm_smfmac_xdlops.hpp:78
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:254
ThisThreadBlock< BlockSize > ThisThreadBlock
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:267
__host__ static constexpr __device__ auto GetABlockDescriptor_Batch_K0PerBlock_MPerBlock_K1()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:343
__host__ static constexpr __device__ auto GetBBlockDescriptor_K0PerBlock_NPerBlock_K1()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:399
__host__ static constexpr __device__ auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:297
__host__ static constexpr __device__ bool CalculateHasMainK0BlockLoop(index_t K0)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:588
ComputeTypeA FloatAAdjusted
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:283
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 AGridDesc_B_K0_M_K1 &a_b_k0_m_k1_grid_desc, const BGridDesc_B_K0_N_K1 &b_b_k0_n_k1_grid_desc, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &c_grid_desc_mblock_mperblock_nblock_nperblock, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CElementwiseOperation &c_element_op, const CBlockClusterAdaptor &c_block_cluster_adaptor)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:641
__host__ static constexpr __device__ auto MakeCBlockClusterAdaptor(const CMNGridDesc &c_m_n_grid_desc, index_t M01, index_t N01, index_t KBatch)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:616
ComputeTypeB FloatBAdjusted
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:284
decltype(MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(CMNGridDesc{})) CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:637
__host__ static constexpr __device__ auto GetBBlockDescriptor_Batch_K0PerBlock_NPerBlock_K1()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:445
decltype(MakeCBlockClusterAdaptor(CMNGridDesc{}, 1, 1, 1)) CBlockClusterAdaptor
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:638
__host__ static constexpr __device__ auto GetCBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:624
static __device__ constexpr bool IsValidCompilationParameter()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:528
__host__ static constexpr __device__ auto MakeCGridDesc_MBlock_MPerBlock_NBlock_NPerBlock(const CMNGridDesc &c_m_n_grid_desc)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:599
__host__ static constexpr __device__ bool CheckValidity(const AGridDesc_B_K0_M_K1 &a_b_k0_m_k1_grid_desc, const BGridDesc_B_K0_N_K1 &b_b_k0_n_k1_grid_desc, const CMNGridDesc &c_m_n_grid_desc, const Block2CTileMap &block_2_ctile_map)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:544
remove_cvref_t< decltype(GridwiseGemmPipeline_Selector< PipelineVer, NumGemmKPrefetchStage >())> GridwiseGemmPipe
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:270
__host__ static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:501
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:25
__host__ constexpr __device__ Merge_v4_no_carry(const LowLengths &low_lengths)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:43
LowLengthsScan low_lengths_scan_
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:38
__host__ constexpr __device__ Merge_v4_no_carry()=default
decltype(make_tuple(container_reduce(LowLengths{}, math::multiplies{}, Number< 1 >{}))) UpLengths
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:35
__host__ static constexpr __device__ bool IsValidUpperIndexMappedToValidLowerIndex(const UpIdx &)
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:116
__host__ __device__ void UpdateLowerIndex(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_up_diff, LowIdx &idx_low, const UpIdx &idx_up_new, Number< Hack >) const
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:81
static constexpr index_t NDimLow
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:26
__host__ static constexpr __device__ index_t GetNumOfLowerDimension()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:52
__host__ constexpr __device__ const auto & GetUpperLengths() const
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:56
__host__ constexpr __device__ void CalculateLowerIndex(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:59
__host__ static constexpr __device__ bool IsKnownAtCompileTime()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:107
__host__ static constexpr __device__ bool IsLinearTransform()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:100
UpLengths up_lengths_
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:39
decltype(container_reverse_exclusive_scan(LowLengths{}, math::multiplies{}, Number< 1 >{})) LowLengthsScan
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:32
__host__ static constexpr __device__ index_t GetNumOfUpperDimension()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:54
__host__ static constexpr __device__ bool IsValidUpperIndexAlwaysMappedToValidLowerIndex()
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:102
LowLengths low_lengths_
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:37
__host__ __device__ void Print() const
Definition: gridwise_gemm_xdlops_bwd_weight.hpp:121
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: thread_group_tensor_slice_transfer_v6r1.hpp:34
Definition: threadwise_tensor_slice_transfer.hpp:39
Definition: integral_constant.hpp:20
Definition: is_known_at_compile_time.hpp:14
Definition: functional2.hpp:33
Definition: unary_element_wise_operation.hpp:334