bar.warp.sync GPU Virtual ISA NVIDIA

bar.warp.sync Parallel Synchronization and Communication Instructions

bar.warp.sync membermask;

bar.warp.sync will cause executing thread to wait until all threads corresponding to membermask have executed a bar.warp.sync with the same membermask value before resuming execution.

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 6.0
Minimum Target sm_30

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.warp.sync membermask; sm_30 bar.warp.sync will cause executing thread to wait until all threads corresponding to membermask have executed a bar.warp.sync with the same membermask value before resuming execution. (see the official PTX ISA docs for the full description)

Operands

  • membermask
    Operand

At a Glance

Data Types -

Related

More in Parallel Synchronization and Communication Instructions

Reference

NVIDIA PTX ISA

Description

bar.warp.sync will cause executing thread to wait until all threads corresponding to membermask have executed a bar.warp.sync with the same membermask value before resuming execution. Operand membermask specifies a 32-bit integer which is a mask indicating threads participating in barrier where the bit position corresponds to thread’s laneid. The behavior of bar.warp.sync is undefined if the executing thread is not in the membermask. bar.warp.sync also guarantee memory ordering among threads participating in barrier. (see the official PTX ISA docs for the full description)

Examples

st.shared.u32 [r0],r1;         // write my result to shared memory
bar.warp.sync  0xffffffff;     // arrive, wait for others to arrive
ld.shared.u32 r2,[r3];         // read results written by other threads

Reproduced from NVIDIA's official PTX ISA documentation for technical accuracy.

Sources