alloca GPU Virtual ISA NVIDIA

alloca Stack Manipulation Instructions

alloca.type ptr, size{, immAlign};

The alloca instruction dynamically allocates memory on the stack frame of the current function and updates the stack pointer accordingly.

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
alloca.type ptr, size{, immAlign}; sm_52 The alloca instruction dynamically allocates memory on the stack frame of the current function and updates the stack pointer accordingly. (see the official PTX ISA docs for the full description)

Operands

At a Glance

Data Types -

Related

More in Stack Manipulation Instructions

Reference

NVIDIA PTX ISA

Description

The alloca instruction dynamically allocates memory on the stack frame of the current function and updates the stack pointer accordingly. The returned pointer ptr points to local memory and can be used in the address operand of ld.local and st.local instructions. If sufficient memory is unavailable for allocation on the stack, then execution of alloca may result in stack overflow. In such cases, attempting to access the allocated memory with ptr will result in undefined program behavior. The memory allocated by alloca is deallocated in the following ways: It is automatically deallocated when the function exits. (see the official PTX ISA docs for the full description)

Semantics

alloca.type ptr, size, immAlign: a = max(immAlign, frame_align); // frame_align is the minimum guaranteed alignment // Allocate size bytes of stack memory with alignment a and update the stack pointer. // Since the stack grows down, the updated stack pointer contains a lower address. stackptr = alloc_stack_mem(size, a); // Return the new value of stack pointer as ptr. Since ptr is the lowest address of the memory // allocated by alloca, the memory can be accessed using ptr up to (ptr + size of allocated memory). stacksave ptr;

Examples

.reg .u32 ra, stackptr, ptr, size;

stacksave.u32 stackptr;     // Save the current stack pointer
alloca ptr, size, 8;        // Allocate stack memory
st.local.u32 [ptr], ra;     // Use the allocated stack memory
stackrestore.u32 stackptr;  // Deallocate memory by restoring the stack pointer

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

Sources