[Question] How is memory ordering guaranteed within an intranode?

In the code, I noticed that the send side uses PTX's 'release' and the receive side uses 'acquire' to ensure memory ordering. However, according to the PTX documentation, 'release' and 'acquire' only guarantee memory ordering within the calling thread. How is memory ordering ensured for other threads within the same warp?

sende：
__device__ __forceinline__ void st_release_sys_global(const int *ptr, int val) {
asm volatile("st.release.sys.global.s32 [%0], %1;"::"l"(ptr), "r"(val) : "memory");
}

recv：
__device__ __forceinline__ uint64_t ld_acquire_sys_global(const uint64_t *ptr) {
uint64_t ret;
asm volatile("ld.acquire.sys.global.u64 %0, [%1];" : "=l"(ret) : "l"(ptr));
return ret;
}

Thank you for your time.

Start with the PTX documentation for the release and acquire operations, then inspect the send and receive helper snippets in the issue. A useful resolution would explain how ordering is ensured for other threads within the same warp, or clarify that additional synchronization is required.

Written by the indexing model from the issue text.