barrier GPU Virtual ISA NVIDIA

Barrier Synchronization Synchronization

bar.sync a{, b};

Block threads in a CTA at a named barrier until the expected number of threads has arrived.

Encoding

PTX is a virtual instruction set. It has no single, stable native binary encoding — the compiler lowers this instruction to different native machine code depending on the selected NVIDIA target architecture (compute capability). This page intentionally shows no bit-diagram; see the target/version requirements below for what governs how this instruction compiles.
PTX ISA Version Introduced PTX ISA 1.0
Minimum Target sm_10

Syntax Forms

One mnemonic covers many type / state-space / scope / modifier combinations — each row below is an independently valid form.

Syntax Data Types State Space(s) Modifiers Min. Target Description
bar.sync a{, b}; sm_10 Wait at barrier resource a for all (or b) threads of the CTA.
barrier.sync a{, b}; sm_10 Alias/expanded form of bar.sync with the same semantics.

Operands

  • a
    Barrier resource identifier
  • b
    Optional thread count participating in the barrier

At a Glance

Data Types —

Related AMDGPU Concepts

Workgroup Barrier ↗
equivalent with restrictions
s_barrier (AMDGPU)

Reference

NVIDIA PTX ISA

Semantics

Every participating thread blocks until all have executed the barrier; also orders shared-memory visibility across the participating threads.

Sources