ld GPU Virtual ISA NVIDIA

Load Memory

ld.space.type d, [a];

Load a value from the specified state space into a register.

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 1.0
Minimum Target sm_10

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.space.type d, [a]; b8, b16, b32, b64, s8, s16, s32, s64, u8, u16, u32, u64, f16, f32, f64 global, local, shared, param, const sm_10 Load from an explicit state space.
ld.global.nc.type d, [a]; global nc sm_35 Load through the read-only (non-coherent) data cache.

Operands

  • d
    Destination register
  • a
    Source address

At a Glance

Data Types b16, b32, b64, b8, f16, f32, f64, s16, s32, s64, s8, u16, u32, u64, u8
State Spaces const, global, local, param, shared
Modifiers nc

Related AMDGPU Concepts

Global-Memory Load ↗
equivalent with restrictions
flat_load_dword (AMDGPU)
Shared/Local-Memory Load ↗
equivalent with restrictions
ds_read_b32 (AMDGPU)

Reference

NVIDIA PTX ISA

Semantics

d = *a, from the given state space.

Sources