Workgroup Barrier Synchronization
Block all threads in a workgroup/CTA at this point until every thread has arrived.
Vendor-Neutral Definition
Every thread in the workgroup/cooperative-thread-array blocks until all participating threads have executed the barrier; may additionally order visibility of shared-memory accesses depending on the exact instruction.
⚠️ Cross-vendor mappings describe semantic relationships and are not guaranteed one-to-one compiler
translations.
equivalent with restrictions
Both block all threads in a workgroup/CTA until every participating thread has arrived.
| Property | NVIDIA PTX | AMDGPU |
|---|---|---|
| Instruction | barrier | s_barrier |
| ISA Layer | Virtual | Native |
| Data Types | — | — |
| Version / Target Introduced | PTX ISA 1.0 | gfx942/gfx1100 |
Important Differences
- PTX bar.sync additionally supports partial-CTA barriers via a thread-count operand and named barrier resources; AMDGPU's s_barrier synchronizes the whole workgroup with no partial-participation form.
Source Evidence
- Parallel Thread Execution ISA ↗ — NVIDIA Corporation
- User Guide for AMDGPU Backend ↗ — LLVM Project
Verification method: documentation