stackrestore GPU Virtual ISA NVIDIA

stackrestore Stack Manipulation Instructions

stackrestore.type a;

Sets the current stack pointer to source register a.

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 7.3
Minimum Target sm_52

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
stackrestore.type a; sm_52 Sets the current stack pointer to source register a. (see the official PTX ISA docs for the full description)

Operands

  • a
    Source operand

At a Glance

Data Types -

Related

More in Stack Manipulation Instructions

Reference

NVIDIA PTX ISA

Description

Sets the current stack pointer to source register a. When stackrestore is used with operand a written by a prior stacksave instruction, it will effectively restore the state of stack as it was before stacksave was executed. Note that if stackrestore is used with an arbitrary value of a, it may cause corruption of stack pointer. This implies that the correct use of this feature requires that stackrestore.type a is used after stacksave.type a without redefining the value of a between them. Operand a has the same type as the instruction type.

Semantics

stackptr = a;

Examples

.reg .u32 ra;
stacksave.u32 ra;
// Code that may modify stack pointer
...
stackrestore.u32 ra;

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

Sources