multimem.cp.async.bulk GPU Virtual ISA NVIDIA
multimem.cp.async.bulk Data Movement and Conversion Instructions
multimem.cp.async.bulk{.sem}.dst.src.completion_mechanism{.cp_mask}
Instruction multimem.cp.async.bulk initiates an asynchronous bulk-copy operation from source address range [srcMem, srcMem + size) to memory locations residing on each GPU’s memory referred to by the destination multimem address range [dstMem, dstMem + size).
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 |
|---|---|---|---|---|---|
| multimem.cp.async.bulk{.sem}.dst.src.completion_mechanism{.cp_mask} | sm_90 | Instruction multimem.cp.async.bulk initiates an asynchronous bulk-copy operation from source address range [srcMem, srcMem + size) to memory locations residing on each GPU’s memory referred to by the… (see the official PTX ISA docs for the full description) |
Operands
At a Glance
Related
Reference
NVIDIA PTX ISA
Description
Instruction multimem.cp.async.bulk initiates an asynchronous bulk-copy operation from source
address range [srcMem, srcMem + size) to memory locations residing on each GPU’s memory referred
to by the destination multimem address range [dstMem, dstMem + size). The direction of
bulk-copy is from the state space specified by the.src modifier to the state space specified
by the.dst modifiers.
The 32-bit operand size specifies the amount of memory to be copied, in terms of number of
bytes. Operand size must be a multiple of 16. The memory range [dstMem, dstMem + size) must not overflow the destination multimem memory space. (see the official PTX ISA docs for the full description)
Examples
multimem.cp.async.bulk.global.shared::cta.bulk_group [dstMem], [srcMem], size;
multimem.cp.async.bulk.global.shared::cta.bulk_group [dstMem], [srcMem], 512;
multimem.cp.async.bulk.global.shared::cta.bulk_group.cp_mask [dstMem], [srcMem], size, byteMask;
// (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.