Fix cute gemm dispatch-5 when C != D

[sjfeng1999]

No description provided.

[thakkarV]

I think this is not a bug. cute::gemm() does not have a way of telling what the first MMA of you sequence is. Whatever collective invokes this gemm is responsible for knowing what the first MMA in the sequence of all tiled MMAs is and zeroing out the D registers before the first MMA. At the scope of a single gemm() dispatch, we cannot know whether k==0 is the first MMA of the sequence or not.

[sjfeng1999]

As I understand it, cute::gemm() computes a gemm with static problem shape (V, M, N, K).
$$D_{mn} = \sum_k A_{mk} * B_{nk} + C_{mn} $$
It is irrelevant with whether it is part of an outer gemm loop or not and whether C and D are zero-initialized.
For example, if we have K = 2, the origin code

CUTE_UNROLL
  for (int k = 0; k < K; ++k) {
    gemm(mma, D, A(_,_,k), B(_,_,k), C);
}

will be unrolled to

gemm(D, A(_, _, 0), B(_, _, 0), C); // Dmn = Am0 x Bn0 + Cmn 
gemm(D, A(_, _, 1), B(_, _, 1), C); // Dmn = Am1 x Bn1 + Cmn 

the result of K=1 will overwrite the result of K=0. This is only true when C = D which is what we do in the CollectiveMainloop (Accum = Amk x Bnk + Accum). But it is incorrect in general case.

Look forward your reply, thks.

[sjfeng1999]

@thakkarV

[github-actions]

This PR has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this PR if it is no longer required. Otherwise, please respond with a comment indicating any updates. This PR will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions]

This PR has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this PR if it is no longer required. Otherwise, please respond with a comment indicating any updates.