wgmma.commit_group GPU Virtual ISA NVIDIA
wgmma.commit_group Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions
wgmma.commit_group.sync.aligned;
wgmma.commit_group instruction creates a new wgmma-group per warpgroup and batches all prior wgmma.mma_async instructions initiated by the executing warp but not committed to any wgmma-group into the new wgmma-group.
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 |
|---|---|---|---|---|---|
| wgmma.commit_group.sync.aligned; | sm_90a | wgmma.commit_group instruction creates a new wgmma-group per warpgroup and batches all prior wgmma. (see the official PTX ISA docs for the full description) |
Operands
At a Glance
Reference
NVIDIA PTX ISA
Description
wgmma.commit_group instruction creates a new wgmma-group per warpgroup and batches all prior wgmma.mma_async instructions initiated by the executing warp but not committed to any
wgmma-group into the new wgmma-group. If there are no uncommitted wgmma.mma_async instructions
then wgmma.commit_group results in an empty wgmma-group.
An executing thread can wait for the completion of all wgmma.mma_async operations in a
wgmma-group by using wgmma.wait_group.
The mandatory.sync qualifier indicates that wgmma.commit_group instruction causes the
executing thread to wait until all threads in the warp execute the same wgmma.commit_group instruction before resuming execution.
The mandatory. (see the official PTX ISA docs for the full description)
Examples
wgmma.commit_group.sync.aligned;Reproduced from NVIDIA's official PTX ISA documentation for technical accuracy.
Sources
- Parallel Thread Execution ISA ↗ - NVIDIA Corporation, Chapter 9 - Instruction Set