[BUG] incorrect spin-loop synchronization

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug
The following sample is using a spin-loop to acquire system-scope writes, but is using relaxed instead of acquire:

cutlass/examples/python/CuTeDSL/distributed/distributed_gemm_reduce_scatter_blackwell.py

Line 1386 in 982748a

res = cute.arch.load(flag.llvm_ptr, cutlass.Int32, sem="relaxed", scope="gpu")

It then uses a warp barrier, but that warp barrier only orders memory operations issued by the same warp, not other warps or other thread blocks, like the example is relying on.

Expected behavior
I expected the spin-loop to use acquire memory ordering.

Additional context
This is a customer request (ping me internally for more details if needed).

@shubaoyu2

@shubaoyu2

thanks for debugging this! here is the fix MR #3118

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.

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.

please ping Bangyu / Xiao on cutedsl distributed example. there is a reason we use relaxed here

both acquire or relaxed work on this case as subsequent read is multimem.ld_reduce, which bypass l1. what acquire does is to CCTL.IVALL invalidate l1. l1 does not play in this story.

for good perf in this case, yes use relax is enough

for general purpose programing, use acquire is a good idea

will have some compiler optimization in future that allow use to use acquire and generate relaxed code based on use case

@XiaoSongXS

Hi @XiaoSongXS, I'm also studying this GEMM-RS kernel. There are 2 additional issues I want to get clarifications from you:

This is using gpu scope for both flag increment and flag polling. Why is sys scope not used here?
There is no fence between relaxed flag polling and LDMC. Shouldn't fence.proxy.alias be necessary? Is it the same reason as "CCTL.IVALL invalidate l1. l1 does not play in this story."? I noticed fence.proxy.alias compiles to MEMBAR.SC.SYS + CCTL.IVALL

This is using gpu scope for both flag increment and flag polling. Why is sys scope not used here?

U need to understand why u need sys scope. if u write data that will flush to ur local gpu l2, then u are using a gpu scope fence (regardless of wether it is ur local gpu thread or remote gpu thread consume this data, as long as the data land on local gpu l2). if u write data that will flush to remote gpu l2, then u need to use sys scope

in this case, the multimem.ld_reduce require all rank have the data in their gmem/l2, which means in epi we write data to local gpu l2, not remote gpu l2, thus the scope of gpu is what we need

There is no fence between relaxed flag polling and LDMC. Shouldn't fence.proxy.alias be necessary? Is it the same reason as "CCTL.IVALL invalidate l1. l1 does not play in this story."? I noticed fence.proxy.alias compiles to MEMBAR.SC.SYS + CCTL.IVALL

no, not the same reason. the principle is the same, u need to think about why u need alias fence. alias fence is needed when u have one PA that map to two VA, and u read/write one VA and later read/write that VA, and u want to ensure the prior VA write is flush to their desired scope (gpu/sys) and is sequential consistency is enforce.

but here, we already have a flag that ensure our data is in l2, and have a spin loop to ensure ldmc happen only after unlock, all those requirement have already satisfied, and thus no alias fence is needed.

alias fence is the most heavy fence, u should try avoid use it. think about what ur program actually need, and how other logic can be use to satisfy what u need.

glad u bring this up and actually do some SASS level analysis, ur heading the right direction on advance cuda programing

U need to understand why u need sys scope. if u write data that will flush to ur local gpu l2, then u are using a gpu scope fence (regardless of wether it is ur local gpu thread or remote gpu thread consume this data, as long as the data land on local gpu l2). if u write data that will flush to remote gpu l2, then u need to use sys scope

Thanks for your insights. The description explains the scope modifier from the perspective of cache flushing behavior. In this way, does that mean it makes no sense to use release.sys for a store in any case as long as the dst is local gmem, since release.gpu has fulfilled the flushing purpose?

And similarly, if I am using regular ld from a local gmem, does that mean acquire.gpu is sufficient, regardless whether the producer is self GPU or remote GPU?

In this way, does that mean it makes no sense to use release.sys for a store in any case as long as the dst is local gmem, since release.gpu has fulfilled the flushing purpose?

Yes. What really matter here is NOT the scope of consumer and producer thread, but rather the scope of where ur memory write goes to.

And the release is only use when u need this "flush" behavior, if u just want local gpu l2 atomic, relaxed is what u need

And similarly, if I am using regular ld from a local gmem, does that mean acquire.gpu is sufficient, regardless whether the producer is self GPU or remote GPU?

Yes, as long as u spin loop pooling from local gpu flag and data, then .acquire.gpu is what u need. The symmetric paring of sys <-> sys is kind of overkill. sys <-> gpu is the minimal hw requirement

@bobboli

@bobboli I see you are a DevTech engineer. Based on Xiao Song's explanation here, I believe many comms kernels in FlashInfer can be further optimized by removing certain defensive fences!

@bobboli I see you are a DevTech engineer. Based on Xiao Song's explanation here, I believe many comms kernels in FlashInfer can be further optimized by removing certain defensive fences!

Unfortunately the A2A kernel in FI uses push to remote & pull from remote pattern, which can not benefit from this. But some others, such as AR, may be optimized.

push to remote can adopt

lamport
counted write
st.async.release the async varions

@XiaoSongXS

@XiaoSongXS dumb question: why are you not worried about the compiler (ptxas) hoisting/reordering multimem.ld_reduce before the relaxed flag poll? My guesses:

there is a while spin control loop going on; or
there is a bar sync in between relaxed load and ld_reduce

But not sure which one is right, or if it is due to other reasons. The SASS code right now shows no reordering was done, but not sure if it is guaranteed to not reorder

this is a excellent question. And ur intuition is exactly correct

when a while loop exist, reorder does not happen.
bar.sync kind of ensure some reorder does not happen, but it is weaker, a few thing it ensure
1. if later operation is cuda core operation and have no memory dependency, then future operation can continue execute. this is what sass defer blocking means. (i think there are lots of hacky sass analysis repo, if ur interested in this, highly recommend dig into the sass)
2. if operation before and after is memory operation, then no reorder will happen (in most case yes, only a few case no). If this is not guaranteed, then below common pattern will break
```
 __shared__ uint32_t shared_scratch_smem;
 shared_scratch_smem[threadId.x] = threadIdx.x + 42;
 __syncthreads(); // bar.sync, need to ensure prior memory op is not reorder after bar.sync, need to ensure later memory op is not reorder prior bar.sync
uint32_t peer_val = shared_scratch_smem[32 - threadIdx.x];
```

A second thought on #3118, yes we should adopt this fix (i'll port this fix in another PR), relies on the compiler to do the optimization, and avoid hacking through the mem sys just to avoid one CCTL.IVALL (little perf hurt, nearly no influence) as it cause customer to be quite confused about what should be the correct model

And one thing on the acquire vs relaxed, we have two concept here that need to be clearly distinguish

wether compiler will reorganize the SASS code that break the release -> acquire pair
wether memory visibility is ensured if compiler dont reorder SASS.

For (1) I think the correct approach is to use .release + .acquire. try not hack through the compiler while loop, as the compiler while loop and bar.sync is "implicit", likely one day it break. the .acquire is a strong ensure provided by the ptx model

For (2) use correct scope for release + acquire will ensure this

@XiaoSongXS

Hey @XiaoSongXS, could I ask about the choice of gpu scope specifically for the flag increment as well?

My understanding of the kernel is that for some of the completed GEMM tiles, the incremented flag should go to the remote peer's L2 rather than a local L2 (since the RS warps use gpu scope polling), so using sys scope is more appropriate. However, I've not yet run into correctness issues from using a gpu scope flag increment.

Is there something else I've failed to consider when reasoning about the right scope for the flag increment?

29. Hey @XiaoSongXS, could I ask about the choice of gpu scope specifically for the flag increment as well?
My understanding of the kernel is that for some of the completed GEMM tiles, the incremented flag should go to the remote peer's L2 rather than a local L2 (since the RS warps use gpu scope polling), so using sys scope is more appropriate. However, I've not yet run into correctness issues from using a gpu scope flag increment.
Is there something else I've failed to consider when reasoning about the right scope for the flag increment?

this is indeed a very interesting question. the short answer is the flag update with release come with two part, (1) a atomic flag update, (2) a member fence. for the atomic flag update, using gpu scope or sys scope in reality have somewhat same hw path and behavior, so that is not the issue, for the member fence, that is the critical path and u need to make sure that part is correct.

so likely u may want to have flag update using gpu scope (as flag is on local gpu l2) and the member with sys scope. for this case u can use the st.release.sys as the sys scope flag update and gpu scope flag update are effectively doing somewhat same thing on hw.

Navigation Menu

Repository navigation

[BUG] incorrect spin-loop synchronization #3117

Which component has the problem?

Bug Report

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

[BUG] incorrect spin-loop synchronization #3117

Description

Which component has the problem?

Bug Report

Activity

brandon-yujie-sun commented on Mar 20, 2026

shubaoyu2 commented on Mar 20, 2026

github-actions commented on Apr 19, 2026

github-actions commented on Jul 18, 2026

XiaoSongXS commented on Jul 20, 2026

XiaoSongXS commented on Jul 20, 2026

gau-nernst commented on Aug 4, 2026

XiaoSongXS commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

bobboli commented on Aug 5, 2026

bobboli commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

gau-nernst commented on Aug 5, 2026

bobboli commented on Aug 5, 2026

XiaoSongXS commented on Aug 5, 2026

henrylhtsang commented on Aug 11, 2026

XiaoSongXS commented on Aug 11, 2026

XiaoSongXS commented on Aug 11, 2026

XiaoSongXS commented on Aug 11, 2026

ShawnWeiChew commented on Oct 26, 2026

XiaoSongXS commented on Oct 27, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions