mbarrier.test_wait GPU Virtual ISA NVIDIA

mbarrier.test_wait Parallel Synchronization and Communication Instructions

// without parity
mbarrier.test_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state;

The test_wait and try_wait operations test for the completion of the current or the immediately preceding phase of an mbarrier object at the location specified by the operand addr.

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 7.0
Minimum Target sm_80

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
// without parity mbarrier.test_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state; sm_80 The test_wait and try_wait operations test for the completion of the current or the immediately preceding phase of an mbarrier object at the location specified by the operand addr. mbarrier. (see the official PTX ISA docs for the full description)

Operands

  • without parity
    Operand

At a Glance

Data Types -

Related

More in Parallel Synchronization and Communication Instructions

Reference

NVIDIA PTX ISA

Description

The test_wait and try_wait operations test for the completion of the current or the immediately preceding phase of an mbarrier object at the location specified by the operand addr. mbarrier.test_wait is a non-blocking instruction which tests for the completion of the phase. mbarrier.try_wait is a potentially blocking instruction which tests for the completion of the phase. If the phase is not complete, the executing thread may be suspended. Suspended thread resumes execution when the specified phase completes OR before the phase completes following a system-dependent time limit. (see the official PTX ISA docs for the full description)

Examples

// Example 1a, thread synchronization with test_wait:

.reg .b64 %r1;
.shared .b64 shMem;

mbarrier.init.shared.b64 [shMem], N;  // N threads participating in the mbarrier.
// (truncated - see the official PTX ISA docs for the full example)

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

Sources