prefetch GPU Virtual ISA NVIDIA

prefetch Data Movement and Conversion Instructions

prefetch{.space}.level [a]; // prefetch to data cache

The prefetch instruction brings the cache line containing the specified address in global or local memory state space into the specified cache level.

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 2.0
Minimum Target sm_20

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
prefetch{.space}.level [a]; // prefetch to data cache sm_20 The prefetch instruction brings the cache line containing the specified address in global or local memory state space into the specified cache level. If the. (see the official PTX ISA docs for the full description)

Operands

At a Glance

Data Types -

Related AMDGPU Concepts

Cache Prefetch ↗
equivalent with restrictions
flat_prefetch_b8 (AMDGPU)

Related

More in Data Movement and Conversion Instructions

Reference

NVIDIA PTX ISA

Description

The prefetch instruction brings the cache line containing the specified address in global or local memory state space into the specified cache level. If the.tensormap qualifier is specified then the prefetch instruction brings the cache line containing the specified address in the.const or.param memory state space for subsequent use by the cp.async.bulk.tensor instruction. If no state space is given, the prefetch uses Generic Addressing. Optionally, the eviction priority to be applied on the prefetched cache line can be specified by the modifier.level::eviction_priority. (see the official PTX ISA docs for the full description)

Examples

prefetch.global.L1             [ptr];
prefetch.global.L2::evict_last [ptr];
prefetchu.L1  [addr];
prefetch.const.tensormap       [ptr];

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

Sources