fabric.try_get GPU Virtual ISA NVIDIA

fabric.try_get Fabric Instructions

fabric.try_get.async.dst.completion_mechanism.sem.scope.b128 [dst], [srcLeId, srcDataOff], size, [bar];

Asynchronously copies size bytes from fabric handle [srcLeId, srcDataOff] to destination memory [dst], where srcLeId is a 32-bit unsigned value denoting the logical endpoint identifier, and…

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.3
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
fabric.try_get.async.dst.completion_mechanism.sem.scope.b128 [dst], [srcLeId, srcDataOff], size, [bar]; sm_100 Asynchronously copies size bytes from fabric handle [srcLeId, srcDataOff] to destination memory [dst], where srcLeId is a 32-bit unsigned value denoting the logical endpoint identifier, and srcDataOff… (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

Asynchronously copies size bytes from fabric handle [srcLeId, srcDataOff] to destination memory [dst], where srcLeId is a 32-bit unsigned value denoting the logical endpoint identifier, and srcDataOff is a 64-bit unsigned value denoting the base offset of the resource to access within the logical endpoint associated with srcLeId. The logical endpoint associated with srcLeId must be a unicast logical endpoint. The size operand is 32 bits and specifies the number of bytes to be copied. It must be a multiple of 16; otherwise, the behavior is undefined. The range [dst, dst + size - 1] must be in bounds of the destination memory space. (see the official PTX ISA docs for the full description)

Examples

fabric.try_get.async.shared::cta.mbarrier::complete_tx::bytes.mbarrier::report::fabric.relaxed.sys.b128 [dstSmem], [srcLeId, srcLeOff], 0x100, [mbar];

fabric.try_get.async.shared::cta.relaxed.sys.mbarrier::complete_tx::bytes.mbarrier::report::fabric.b128 [dstSmem], [srcLeId, srcLeOff], sizeBytes, [mbar];

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

Sources