st GPU Virtual ISA NVIDIA

Store Data Movement and Conversion Instructions

st.space.type [a], b;

Store a register value into the specified state space.

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
st.space.type [a], b; b8, b16, b32, b64, s8, s16, s32, s64, u8, u16, u32, u64, f16, f32, f64 global, local, shared, param sm_10 Store to an explicit state space.

Operands

  • a
    Destination address
  • b
    Value to store

At a Glance

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

Related AMDGPU Concepts

Global-Memory Store ↗
equivalent with restrictions
Shared/Local-Memory Store ↗
equivalent with restrictions
ds_write_b32 (AMDGPU)
Private and Scratch Memory Access ↗
equivalent with restrictions

Related

More in Data Movement and Conversion Instructions

Reference

NVIDIA PTX ISA

Description

Store the value of operand b in the location specified by the destination address operand a in specified state space. If no state space is given, perform the store using Generic Addressing. Stores to const memory are illegal. If no sub-qualifier is specified with.shared state space, then::cta is assumed by default. Supported addressing modes for operand a and alignment requirements are described in Addresses as Operands. If.param is specified without any sub-qualifiers then it defaults to.param::func. Instruction st.param{::func} used for passing arguments to device function cannot be predicated. (see the official PTX ISA docs for the full description)

Semantics

*a = b, in the given state space.

Examples

st.global.f32    [a],b;
st.local.b32     [q+4],a;
st.global.v4.s32 [p],Q;
st.local.b32     [q+-8],a; // negative offset
st.local.s32     [100],r7; // immediate address
// (truncated - see the official PTX ISA docs for the full example)

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

Sources