jinhongyii opened a new pull request, #20074:
URL: https://github.com/apache/tvm/pull/20074
## Motivation and context
`MBarrier.remote_view(rank)` represents an mbarrier owned by another CTA in
the same cluster. The existing view kept only a buffer whose pointer had been
mapped to the remote CTA with PTX `mapa`. Calling `remote_bar.arrive(...)`
then followed the inherited local-arrive path and emitted the local
`mbarrier.arrive.shared.b64` form against that mapped remote address.
A remote arrival must instead use the `shared::cluster` instruction form.
TIRx models that form with the owner-local barrier pointer plus the destination
CTA rank and predicate. Using the local instruction with a remote address is
not equivalent and is reported by synccheck as a local arrival on a remote
mbarrier address.
## Changes
- Keep the owner-local buffer and target CTA rank when constructing a remote
mbarrier view.
- Route `MBarrier` remote arrivals through the cluster helper using the
local barrier pointer and stored CTA rank.
- Apply the same routing to `TMABar`, including `mbarrier.arrive.expect_tx`.
- Keep a typed `PointerType(uint64, shared)` mapped buffer on the view so
`ptr_to` remains available to operations that explicitly consume a remote
shared-memory pointer.
- Reject operations with ambiguous or invalid ownership semantics:
- initializing or waiting on a remote view,
- supplying another `cta_id` to a view that already fixes its target,
- creating a remote view from another remote view.
- Preserve the existing local-CTA behavior for ordinary barriers.
## Testing
- Verify the typed `mapa` binding and remote buffer in TIRx IR.
- Verify CUDA codegen for plain and counted
`mbarrier.arrive.shared::cluster.b64` forms.
- Verify CUDA codegen for remote
`mbarrier.arrive.expect_tx.shared::cluster.b64`.
- Verify that the corresponding local instruction forms are not emitted for
remote views.
- Verify diagnostics for remote init, wait, nested views, and conflicting
`cta_id` arguments.
- Run changed-files pre-commit checks.
Focused result: 3 tests passed.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
