[QST]Question about Utccpop

[zhang662817]

https://github.com/NVIDIA/cutlass/blob/c2ad7c5b20f131c4ba33601860f1da3f9c9df0f3/include/cutlass/gemm/collective/sm100_blockscaled_mma_warpspecialized.hpp#L834

For sm_pair case， ScaleFactorB in smem should multicast to leader and peer cta tmem, which triggered by leader cta; ScaleFactorA in leader cta smem should copy to leader tmem and ScaleFactorA in peer cta smem should copy to peer tmem, which both triggered by leader cta.

Is my understanding correct?

If correct, how tcgen05.cp.cta_group::2 meet this reqeust?

Thanks.

[sazczmh]

Based on my understanding, in the CTA-Pair scenario, smem_SFB maintains a complete independent copy per SM. This behavior differs fundamentally from smem_B, which does not possess this per-SM duplication characteristic.

As shown in this example, when the Cluster configuration is set to (2,1,1), only smem_SFB undergoes broadcasting."

When the Cluster configuration is set to (2,2,1), smem_A, smem_SFA, smem_SFB undergoes broadcasting.

Therefore, my understanding is that tcgen05.cp.cta_group::2 lacks the capability for SMEM broadcasting across two CTAs — though it can broadcast to four warps.

[zhang662817]

Ok, thanks.

ScaleFactor in leader cta smem is copied to leader tmem and ScaleFactor in peer cta smem is copied to peer tmem, which both triggered by leader cta.

When .cta_group::2 is specified, the data is copied into the Tensor Memory of both the current and the peer CTAs. But the source is in Individual smem.

[github-actions[bot]]

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.

[github-actions[bot]]

This issue has been labeled inactive-90d due to no recent activity in the past 90 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.