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 systemdependent time limit. The optional 32-bit unsigned integer operand suspendTimeHint specifies the
time limit, in nanoseconds, that may be used for the time limit instead of the system-dependent limit.
mbarrier.test_wait and mbarrier.try_wait test for completion of the phase :
▶ Specified by the operand state, which was returned by an mbarrier.arrive instruction on the
same mbarrier object during the current or the immediately preceding phase. Or
▶ Indicated by the operand phaseParity, which is the integer parity of either the current phase
or the immediately preceding phase of the mbarrier object.
The .parity variant of the instructions test for the completion of the phase indicated by the operand
phaseParity, which is the integer parity of either the current phase or the immediately preceding
phase of the mbarrier object. An even phase has integer parity 0 and an odd phase has integer parity
of 1. So the valid values of phaseParity operand are 0 and 1.
Note: the use of the .parity variants of the instructions requires tracking the phase of an mbarrier
object throughout its lifetime.
How should phase parity be used in mbarrier operations? I have tried arriving without specifying parity (since arrival doesn’t require it), followed by a wait where I must set the parity. If I set it to 0, it freezes; if I set it to 1, it runs fine. In the FA3 code, wait uses (work_idx + 1) % 2, which allows parity to be either 0 or 1. I don’t understand how to select the correct parity.