[QST] Why we use three sync in sgemm_1?

[ziyuhuang123]

I am learning this example:

cutlass/examples/cute/tutorial/sgemm_1.cu

Lines 209 to 211 in c4e3e12

	cp_async_fence(); // Label the end of (potential) cp.async instructions
	cp_async_wait<0>(); // Sync on all (potential) cp.async instructions
	__syncthreads(); // Wait for all threads to write to smem

What is your question?

    cp_async_fence();        // Label the end of (potential) cp.async instructions
    cp_async_wait<0>();      // Sync on all (potential) cp.async instructions
    __syncthreads();         // Wait for all threads to write to smem

Why we have three sync in sgemm_1???

[cloudhan]

On some new architectures say sm_80, the very basic form copy(...), that is, without specifying the copy atom, might generate cp.async instruction. So you need further safety guarantee.

#1231

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.