Understanding sync and arrive in FA3 Store Function

[ziyuhuang123]

Does anyone understand why the number of sync operations in the FA3 store function (

flash-attention/hopper/epilogue_bwd_sm90_tma.hpp

Line 161 in 0dfb281

cutlass::arch::NamedBarrier::arrive(NumEpilogueThreads + cutlass::NumThreadsPerWarp,

) is 256+32 (within a single warp, with all threads arriving)? Why does the non-varlen case require arrive before sync, while the varlen case only needs sync without arrive?