Operands

3. Operands#

PISA instruction has zero or more operands, whose size and type are defined by the instruction template. No implicit type conversion is performed on the operands, and each operand may be a register variable, an immediate value, or an address.

3.1. Register Operands#

A register operand is specified using the name of a register or a predicate variable.

.reg .32b %dst, %src0, %src1;
.pred %pin, %pout;

iadd.32b %dst, %src0, %src1;
uaddc.ci.co.32b %dst, %pout, %src0, %src1, %pin;

3.2. Immediate Values#

An immediate value can be used as a source operand anywhere a scalar register variable is allowed, unless specified otherwise by the instruction template.

An immediate value used in integer instructions can be specified as a decimal or a hexadecimal literal (prefixed with 0x or 0X). Hexadecimal literals are zero-extended, while decimal literals are sign-extended to 64 bits. If the resulting 64-bit value does not mathematically fit within the signed or unsigned range of the operand size, a parser error will be generated.

 .reg .32b %dst, %src0;

 // add -1 to %src0 and store the result in %dst
 iadd.32b %dst, %src0, -1;
// add 15 to %src0 and store the result in %dst
 iadd.32b %dst, %src0, 0xF;

An immediate value used in floating-point instructions may either be a decimal floating-point literal, or in IEEE hexadecimal format. Decimal floating-point literal is treated as a double-precision value, which is converted to type of an operand using implementation-defined rounding and denorm mode.

Important

Size of hexadecimal floating-point value must match size of operand, as no type conversion is allowed.

 .reg .32b %dst, %src0;

 // add single-precision 4.2 to %src0 and store the result in %dst
 fadd.f %dst, %src0, 4.2;
// add single-precision -1.0 to %src0 and store the result in %dst
 fadd.f %dst, %src0, 0xBF800000;

Note

There is no provision for a vector of immediate values.

3.3. Address Operands#

Address operands are used in memory-related instructions and are written in the form [Value1] or [Value1 + Value2]. Instruction descriptions reference the following shorthand names for common operand forms:

Address operand forms#

Name

Meaning

Syntax

MemRR

32-bit or 64-bit base address with a register offset

[reg + reg], [var + reg]

MemRI

32-bit or 64-bit base address with a signed immediate offset

[reg + imm], [var + imm]

MemR

32-bit or 64-bit base address without an explicit offset

[reg], [var]

MemI

Immediate address value without an explicit offset

[imm]

Here reg denotes a scalar register, var denotes a memory variable, and imm denotes an immediate value accepted by the instruction template.

Warning

Not all instructions support all operand forms.

A memory variable represents the byte address of a variable in a memory address space. The address size is interpreted as follows:

  • 32-bit value for .private and .shared variables.

  • 64-bit value for .global and .const variables.

  • 32 or 64-bit value for a register variable, depending on its declaration.

Either 32-bit or 64-bit arithmetic is performed based on the address operand values, and the result serves as the memory access address.

.reg .32b %reg32;
.const .align 4 @constVar[4] = { .32b 42 };
.shared .align 4 @sharedVar[16];

// load value from global address space specified by %reg64
ld.global.32b %reg32, [%reg64];
// load value from 32bit shared address space specified by @sharedVar[4]
ld.shared.32b %reg32, [@sharedVar + 4];
// load value from 64bit const address space specified by @constVar[0]
ld.const.32b %reg32, [@constVar];

Warning

An out-of-bounds memory access results in undefined behavior.

3.3.1. Generic Addressing#

Address operands also support a generic addressing mode where the address space is not specified at compile time. In this mode, the address operand is treated as a 64-bit generic address, with translation to the concrete address space being accomplished via addrcast instruction.

 .private @A[20];
 .reg .64b %gen64;
 .reg .32b %prv32, %reg32;

 // convert 32-bit private address into a 64-bit generic address
 addrcast.generic.private %gen64, @A;
// load 4 bytes from offset 12 of @A
 ld.generic.32b %reg32, [%gen64 + 12];
 // update %gen64 to point to @A[4]
 iadd.64b %gen64, %gen64, 4;
 // convert 64-bit generic address back to 32-bit private address
 addrcast.private.generic %prv32, %gen64;
// load 4 bytes from @A[4] using the converted private address
 ld.private.32b %reg32, [%prv32];

3.4. Vector Operands#

Register operands of vector type are supported by a limited number of memory access and data movement instructions. For all other instructions, each element of the vector must be accessed individually.

Elements of a vector operand can be accessed by specialized instructions (e.g. extract and insert) or via swizzle syntax that allows individual access to the first 4 elements of a vector.

Swizzle Values#

Swizzle

Element

.x

0

.y

1

.z

2

.w

3

 
.reg .v2.32b %vec;
.reg .32b %dst, %src;

// load 2 32-bit values
ld.private.v2.32b %vec, [%addr];
// access %vec[0]
mov.32b %dst, %vec.x;
// add 2 elements of %vec
iadd.32b %dst, %dst, %vec.y;

mov instruction supports additional multi-element swizzles (.xy, .zw, and .xyzw) to enable data movement between vector and scalar variables.

Contents