[FIX][TIRx] Use cluster arrivals for remote mbarrier views

[jinhongyii]

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.

[gemini-code-assist]

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.