fabric.try_atom GPU Virtual ISA NVIDIA

Fabric Instructions

// Bitwise operations
fabric.try_atom.async.src.completion_mechanism{.level::cache_hint}.sem.sco.redOpBit.typeBit
[dstLeId, dstDataOff], [smem_dst], [smem_src], [mbar] {, cache_policy};

Atomically loads the original value at location [dstLeId, dstDataOff] into destination shared memory location [smem_dst], performs a reduction operation with the data at shared memory location…

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 9.4
Minimum Target sm_100

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
// Bitwise operations fabric.try_atom.async.src.completion_mechanism{.level::cache_hint}.sem.sco.redOpBit.typeBit [dstLeId, dstDataOff], [smem_dst], [smem_src], [mbar] {, cache_policy}; sm_100 Atomically loads the original value at location [dstLeId, dstDataOff] into destination shared memory location [smem_dst], performs a reduction operation with the data at shared memory location… (see the official PTX ISA docs for the full description)

Operands

At a Glance

Data Types -

Related

More in Fabric Instructions

Reference

NVIDIA PTX ISA

Description

Atomically loads the original value at location [dstLeId, dstDataOff] into destination shared memory location [smem_dst], performs a reduction operation with the data at shared memory location [smem_src] and the value in location [dstLeId, dstDataOff], and stores the result of the specified operation at location [dstLeId, dstDataOff], overwriting the original value. For the .cas (compare-and-swap) operation, the lower 16 bytes at the location [smem_src] is the compare value and the upper 16 bytes at the location [smem_src] is the swap value. (see the official PTX ISA docs for the full description)

Semantics

atomic { [smem_dst] = [dstLeId, dstDataOff]; // copy as the prev value if (operation == cas) { // smem_src must be 32B aligned b = lower16Bytes([smem_src]); // smem_src value to be compared c = upper16Bytes([smem_src]); // swap value [dstLeId, dstDataOff] = cas([dstLeId, dstDataOff], b, c); } else { [dstLeId, dstDataOff] = operation([dstLeId, dstDataOff], [smem_src]); } } exch(r, s) = s; cas(r, s, t) = (r == s) ? t : r;

Examples

fabric.try_atom.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.and.b32 [dstLeId, dstLeOffData], [dstSmem], [srcSmem], [mbar];

fabric.try_atom.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.add.bf16x2 [dstLeId, dstLeOffData], [dstSmem], [srcSmem], [mbar];

fabric.try_atom.async.shared::cta.L2::cache_hint.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.xor.b64 [dstLeId, dstLeOffData], [dstSmem], [srcSmem], [mbar], mdesc;

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

Sources