ld.global.nc GPU Virtual ISA NVIDIA

ld.global.nc Data Movement and Conversion Instructions

ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache_policy};

Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent read-only cache.

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 3.1
Minimum Target sm_32

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
ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache_policy}; sm_32 Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent read-only cache. (see the official PTX ISA docs for the full description)

Operands

At a Glance

Data Types -

Related

More in Data Movement and Conversion Instructions

Reference

NVIDIA PTX ISA

Description

Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent read-only cache. Note On some architectures, the texture cache is larger, has higher bandwidth, and longer latency than the global memory cache. For applications with sufficient parallelism to cover the longer latency, ld.global.nc should offer better performance than ld.global on such architectures. The address operand a shall contain a global address. Supported addressing modes for operand a and alignment requirements are described in Addresses as Operands. The.v8 (.vec ) qualifier is supported if:.type is.b32,.s32,.u32, or. (see the official PTX ISA docs for the full description)

Semantics

d = a; // named variable a d = *(&a+immOff) // variable-plus-offset d = *a; // register d = *(a+immOff); // register-plus-offset d = *(immAddr); // immediate address

Examples

ld.global.nc.f32           d, [a];
ld.gloal.nc.L1::evict_last.u32 d, [a];

createpolicy.fractional.L2::evict_last.b64 cache_policy, 0.5;
ld.global.nc.L2::cache_hint.f32  d, [a], cache_policy;
// (truncated - see the official PTX ISA docs for the full example)

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

Sources