[QST]synchronize before tma store

[haitianzhishang]

What is your question?
What is your question?
In sm90_epilogue_array_tma_warpspecialized.hpp

Image Is it necessary to add "fence_view_async_shared" before “synchronize”? I think "synchronize" has ensured that tma store will be issued after the movement between register and share memory has completed.
If "fence_view_async_shared" is necessary, does it mean "synchronize" has no effect on tma store?

[CalebDu]

Hi @haitianzhishang Thanks for question. TMA instruction(cp.async.bulk) belongs to async proxy. LDS instruction belongs to generic proxy . You can find more details in PTX doc

Accessing the same memory location across multiple proxies needs a cross-proxy fence. For the async proxy, fence.proxy.async should be used to synchronize memory between generic proxy and the async proxy.

But fence.proxy.async can not sync among threads, only check(sync) single thread memory transaction completed. So there is a synchronize()[namedbarrier] to sync all epilogue threads.

[haitianzhishang]

Thanks for your reply