Parallel Synchronization and Communication Instructions Instructions
31 PTX instructions in this category - click any row for encoding, pseudocode, and full documentation.
| Mnemonic | Syntax | PTX Version | Summary |
|---|---|---|---|
| activemask | activemask.b32 d; | PTX ISA 6.2 | Query the bitmask of currently active (converged) lanes in the executing warp. |
| atom | atom.space.op.type d, [a], b; | PTX ISA 1.1 | Atomically read-modify-write a memory location and return the prior value. |
| bar.cta | bar.cta.sync a{, b}; | PTX ISA 1.0 | Performs barrier synchronization and communication within a CTA. |
| bar.warp.sync | bar.warp.sync membermask; | PTX ISA 6.0 | 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. |
| barrier | barrier{.cta}.sync{.aligned} a{, b}; | PTX ISA 1.0 | Block threads in a CTA at a named barrier until the expected number of threads has arrived. |
| barrier.cluster | barrier.cluster.arrive{.sem}{.aligned}; | PTX ISA 7.8 | Performs barrier synchronization and communication within a cluster. |
| barrier.cta | barrier.cta.sync{.aligned} a{, b}; | PTX ISA 1.0 | Performs barrier synchronization and communication within a CTA. |
| clusterlaunchcontrol.query_cancel | clusterlaunchcontrol.query_cancel.is_canceled.pred.b128 pred, try_cancel_response; | PTX ISA 8.6 | Instruction clusterlaunchcontrol.query_cancel can be used to decode opaque response written by instruction clusterlaunchcontrol.try_cancel. |
| clusterlaunchcontrol.try_cancel | clusterlaunchcontrol.try_cancel.async{.space}.completion_mechanism{.multicast::cluster::all}.b128 [addr], [mbar]; | PTX ISA 8.6 | The clusterlaunchcontrol.try_cancel instruction requests atomically cancelling the launch of a cluster that has not started running yet. |
| cp.async.mbarrier.arrive | cp.async.mbarrier.arrive{.noinc}{.shared{::cta}}.b64 [addr]; | PTX ISA 7.0 | Causes an arrive-on operation to be triggered by the system on the mbarrier object upon the completion of all prior cp.async operations initiated by the executing thread. |
| elect.sync | elect.sync d|p, membermask; | PTX ISA 8.0 | elect.sync elects one predicated active leader thread from among a set of threads specified by membermask. |
| fence | // Thread fence: fence{.sem}.scope; |
PTX ISA 1.4 | The fence instruction establishes an ordering between memory accesses requested by this thread, as described by the memory consistency model. |
| griddepcontrol | griddepcontrol.action; | PTX ISA 7.8 | The griddepcontrol instruction allows the dependent grids and prerequisite grids as defined by the runtime, to control execution in the following way: |
| match.sync | match.any.sync.type d, a, membermask; | PTX ISA 6.0 | match.sync will cause executing thread to wait until all non-exited threads from membermask have executed match.sync with the same qualifiers and same membermask value before resuming execution. |
| mbarrier.arrive | mbarrier.arrive{.sem.scope}{.shared{::cta}}.b64 state, [addr]{, count}; | PTX ISA 7.0 | A thread executing mbarrier.arrive performs an arrive-on operation on the mbarrier object at the location specified by the address operand addr. The 3 |
| mbarrier.arrive_drop | mbarrier.arrive_drop{.sem.scope}{.shared{::cta}}.b64 state, [addr] {, count}; | PTX ISA 7.0 | A thread executing mbarrier.arrive_drop on the mbarrier object at the location specified by the address operand addr performs the following steps: Decrements the expected arrival count of the mbarrier object by the value specified by the 32-bit integer operand count. |
| mbarrier.check_layout | mbarrier.check_layout.layout{.ss}.b64 p, [addr]; | PTX ISA 9.3 | The layout of the opaque mbarrier object can be queried using mbarrier.check_layout. |
| mbarrier.complete_tx | mbarrier.complete_tx{.sem.scope}{.space}.b64 [addr], txCount; | PTX ISA 8.0 | A thread executing mbarrier.complete_tx performs a complete-tx operation on the mbarrier object at the location specified by the address operand addr. |
| mbarrier.expect_tx | mbarrier.expect_tx{.sem.scope}{.space}.b64 [addr], txCount; | PTX ISA 8.0 | A thread executing mbarrier.expect_tx performs an expect-tx operation on the mbarrier object at the location specified by the address operand addr. |
| mbarrier.init | mbarrier.init{.layout}{.shared{::cta}}.b64 [addr], count; | PTX ISA 7.0 | mbarrier.init initializes the mbarrier object at the location specified by the address operand addr with the unsigned 32-bit integer count. |
| mbarrier.inval | mbarrier.inval{.shared{::cta}}.b64 [addr]; | PTX ISA 7.0 | mbarrier.inval invalidates the mbarrier object at the location specified by the address operand addr. |
| mbarrier.pending_count | mbarrier.pending_count{.layout}.b64 count, state; | PTX ISA 7.0 | The pending count can be queried from the opaque mbarrier state using mbarrier.pending_count. |
| mbarrier.test_wait | // without parity mbarrier.test_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state; |
PTX ISA 7.0 | 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.try_wait | // without parity mbarrier.try_wait{.phase_type::primary}{.sem.scope}{.ss}.b64 waitComplete, [addr], state {, timeHint}; |
PTX ISA 7.0 | 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. |
| membar | membar.level; | PTX ISA 1.4 | Order this thread's prior memory accesses relative to later ones, visible to a given scope. |
| multimem.red.async | multimem.red.async.sem.scope{.ss}.op.type [a], b; | PTX ISA 9.3 | multimem.red.async is a non-blocking instruction which initiates an asynchronous reduction operation specified by.op, with operand b and the value at memory locations residing on each GPU’s memory referred to by the destination multimem address operand a. |
| red | red.space.op.type [a], b; | PTX ISA 1.2 | Atomically read-modify-write a memory location without returning the prior value. |
| red.async | // Increment and Decrement reductions red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar]; |
PTX ISA 8.1 | red.async is a non-blocking instruction which initiates an asynchronous reduction operation specified by.op, with the operand b and the value at destination shared memory location specified by operand a. |
| redux.sync | redux.sync.op.type dst, src, membermask; | PTX ISA 7.0 | redux.sync will cause the executing thread to wait until all non-exited threads corresponding to membermask have executed redux.sync with the same qualifiers and same membermask value before resuming execution. |
| tensormap.cp_fenceproxy | tensormap.cp_fenceproxy.cp_qualifiers.fence_qualifiers.sync.aligned [dst], [src], size; | PTX ISA 8.3 | The tensormap.cp_fenceproxy instructions perform the following operations in order: Copies data of size specified by the size argument, in bytes, from the location specified by the address operand… |
| vote | vote.mode.pred d, {!}a; | PTX ISA 6.0 | Combine a per-lane predicate across the warp using any/all/ballot reduction. |