tcgen05.commit GPU Virtual ISA NVIDIA
tcgen05.commit TensorCore 5th Generation Family Instructions
tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64
The instruction tcgen05.commit is an asynchronous instruction which makes the mbarrier object, specified by the address operand mbar, track the completion of all the prior asynchronous tcgen05 operations, as listed in mbarrier based completion mechanism, initiated by the executing thread.
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.
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 |
|---|---|---|---|---|---|
| tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64 | sm_100a | The instruction tcgen05.commit is an asynchronous instruction which makes the mbarrier object, specified by the address operand mbar, track the completion of all the prior asynchronous tcgen05 operati… (see the official PTX ISA docs for the full description) |
Operands
At a Glance
Related
Reference
NVIDIA PTX ISA
Description
The instruction tcgen05.commit is an asynchronous instruction which makes the mbarrier object,
specified by the address operand mbar, track the completion of all the prior asynchronous tcgen05 operations, as listed in mbarrier based completion mechanism,
initiated by the executing thread. Upon the completion of the tracked asynchronous tcgen05 operations, the signal specified by the.completion_mechanism is triggered by the system
on the mbarrier object.
This instruction accesses its mbarrier operand using generic-proxy.
The instruction tcgen05.commit.cta_group::1 tracks for the completion of all prior
asynchronous tcgen05 operations with.cta_group::1 issued by the current thread. (see the official PTX ISA docs for the full description)
Examples
Example 1:
tcgen05.cp.cta_group::1.128x256b [taddr0], sdesc0;
tcgen05.commit.cta_group::1.mbarrier::arrive::one.b64 [mbarObj1];
loop:
mbarrier.try_wait.parity.b64 p, [mbarObj1], 0;
// (truncated - see the official PTX ISA docs for the full example)Reproduced from NVIDIA's official PTX ISA documentation for technical accuracy.
Sources
-
Parallel Thread Execution ISA ↗
- NVIDIA Corporation, Chapter 9 - Instruction Set
Deep-linked directly to this instruction's section.