Syntax and Format

1. Syntax and Format#

1.1. Program#

A PISA program consists of one or more modules, with each module being defined in a separate source file. One or more modules can be combined into a single module as either an executable (with entry kernels) or a library.

A module contains a header, followed by one or more kernel, function, and global variable declarations and definitions. Kernels, functions, and global variables may also have attributes specified alongside them. A module may also include debugging, profiling, and other identifying information.

A global variable is a variable that is declared at the module level and is accessible by all kernels and functions within the module. Global variables are identified via .global or .const directives.

A kernel is a special type of function that serves as an entry point for execution on the device, and can only be invoked from the host program. It is identified via .kernel directive and has no return value.

A function is a subroutine that can be invoked by kernels or other functions. It is identified via .function directive and may have a return value.

Example PISA Module#
// global variable declaration
.global .align 4 @A[4] = { .32b 42 };

// function returning 32-bit value
.function .32b @funcName() {
  .reg .32b %r0;

  ld.global.32b %r0, [@A];
  return %r0;
}

// kernel entry point
.kernel @kernelName() {
  .reg .32b %rv;

  call %rv, @funcName();
}

1.2. Statement#

A PISA kernel or function is a collection of statements. A statement is a variable declaration, instruction, label, scope, or comment.

A declaration introduces a variable and allocates storage for it. A declaration must terminate with a ;.

A label is an identifier that marks a position in the code and is used as a target for goto instruction. Labels must appear alone on a line and terminate with a :.

An instruction specifies an operation to be performed, such as arithmetic, logic, data movement, etc. Instruction must terminate with a ;.

A scope is a block of statements enclosed in curly braces {}. Variables declared within a scope are visible throughout that scope and any nested subscopes.

A comment provides explanatory notes for human readers and is ignored by the assembler. Comments begin with a // and continue to the end of the line.

Example PISA statements#
// variable declaration
.reg .32b %r, %a, %b;

label1:
  // instruction
  iadd.32b   %r, %a, %b;

1.3. Identifiers#

Named identifiers are used for functions, kernels, variables, registers, and labels. A named identifier consists of an optional prefix character and the name string.

The name string must match regular expression [a-zA-Z$._][a-zA-Z$._0-9]*.

  • ‘@’ prefix is used for functions, kernels, and memory variable identifiers.

  • ‘%’ prefix is used for register identifiers.

  • label identifiers do not have a prefix character.

Contents