[GH-ISSUE #380] [Question] How is memory ordering guaranteed within an intranode? #258

Open
opened 2026-10-26 09:57:30 +00:00 by bsakaguchi · 1 comment
Owner

Originally created by @dsadsadedsadsa on GitHub (Aug 28, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/380

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.

Originally created by @dsadsadedsadsa on GitHub (Aug 28, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/380 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.
Author
Owner

@sphish commented on GitHub (Sep 24, 2025):

I suggest you read the section on memory ordering in the PTX documentation. https://docs.nvidia.com/cuda/parallel-thread-execution/#ordering-memory-operations

<!-- gh-comment-id:3326353278 --> @sphish commented on GitHub (Sep 24, 2025): I suggest you read the section on memory ordering in the PTX documentation. https://docs.nvidia.com/cuda/parallel-thread-execution/#ordering-memory-operations
Sign in to join this conversation.
No labels
pull-request
No milestone
No project
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepEP#258
No description provided.