MLIR

Multi-Level IR Compiler Framework

'acc' Dialect

An OpenACC dialect for MLIR. This dialect models the construct from the OpenACC 3.3 directive language.

Operation definition 

acc.atomic.capture (acc::AtomicCaptureOp) 

Performs an atomic capture

Syntax:

operation ::= `acc.atomic.capture` $region attr-dict

This operation performs an atomic capture.

The region has the following allowed forms:

  acc.atomic.capture {
    acc.atomic.update ...
    acc.atomic.read ...
    acc.terminator
  }

  acc.atomic.capture {
    acc.atomic.read ...
    acc.atomic.update ...
    acc.terminator
  }

  acc.atomic.capture {
    acc.atomic.read ...
    acc.atomic.write ...
    acc.terminator
  }

Traits: RecursiveMemoryEffects, SingleBlock, SingleBlockImplicitTerminator

Interfaces: AtomicCaptureOpInterface

acc.atomic.read (acc::AtomicReadOp) 

Performs an atomic read

Syntax:

operation ::= `acc.atomic.read` $v `=` $x
              `:` type($x) `,` $element_type attr-dict

This operation performs an atomic read.

The operand x is the address from where the value is atomically read. The operand v is the address where the value is stored after reading.

Interfaces: AtomicReadOpInterface

Attributes: 

AttributeMLIR TypeDescription
element_type::mlir::TypeAttrany type attribute

Operands: 

OperandDescription
xpointer-like type
vpointer-like type

acc.atomic.update (acc::AtomicUpdateOp) 

Performs an atomic update

Syntax:

operation ::= `acc.atomic.update` $x `:` type($x) $region attr-dict

This operation performs an atomic update.

The operand x is exactly the same as the operand x in the OpenACC Standard (OpenACC 3.3, section 2.12). It is the address of the variable that is being updated. x is atomically read/written.

The region describes how to update the value of x. It takes the value at x as an input and must yield the updated value. Only the update to x is atomic. Generally the region must have only one instruction, but can potentially have more than one instructions too. The update is sematically similar to a compare-exchange loop based atomic update.

The syntax of atomic update operation is different from atomic read and atomic write operations. This is because only the host dialect knows how to appropriately update a value. For example, while generating LLVM IR, if there are no special atomicrmw instructions for the operation-type combination in atomic update, a compare-exchange loop is generated, where the core update operation is directly translated like regular operations by the host dialect. The front-end must handle semantic checks for allowed operations.

Traits: RecursiveMemoryEffects, SingleBlock, SingleBlockImplicitTerminator

Interfaces: AtomicUpdateOpInterface

Operands: 

OperandDescription
xpointer-like type

acc.atomic.write (acc::AtomicWriteOp) 

Performs an atomic write

Syntax:

operation ::= `acc.atomic.write` $x `=` $expr
              `:` type($x) `,` type($expr)
              attr-dict

This operation performs an atomic write.

The operand x is the address to where the expr is atomically written w.r.t. multiple threads. The evaluation of expr need not be atomic w.r.t. the write to address. In general, the type(x) must dereference to type(expr).

Interfaces: AtomicWriteOpInterface

Operands: 

OperandDescription
xpointer-like type
exprany type

acc.attach (acc::AttachOp) 

Represents acc attach semantics which updates a pointer in device memory with the corresponding device address of the pointee.

Syntax:

operation ::= `acc.attach` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.bounds (acc::DataBoundsOp) 

Represents normalized bounds information for acc data clause.

Syntax:

operation ::= `acc.bounds` oilist(
              `lowerbound` `(` $lowerbound `:` type($lowerbound) `)`
              | `upperbound` `(` $upperbound `:` type($upperbound) `)`
              | `extent` `(` $extent `:` type($extent) `)`
              | `stride` `(` $stride `:` type($stride) `)`
              | `startIdx` `(` $startIdx `:` type($startIdx) `)`
              ) attr-dict

This operation is used to record bounds used in acc data clause in a normalized fashion (zero-based). This works well with the PointerLikeType requirement in data clauses - since a lowerbound of 0 means looking at data at the zero offset from pointer.

The operation must have an upperbound or extent (or both are allowed - but not checked for consistency). When the source language’s arrays are not zero-based, the startIdx must specify the zero-position index.

Examples below show copying a slice of 10-element array except first element. Note that the examples use extent in data clause for C++ and upperbound for Fortran (as per 2.7.1). To simplify examples, the constants are used directly in the acc.bounds operands - this is not the syntax of operation.

C++:

int array[10];
#pragma acc copy(array[1:9])

=>

acc.bounds lb(1) ub(9) extent(9) startIdx(0)

Fortran:

integer :: array(1:10)
!$acc copy(array(2:10))

=>

acc.bounds lb(1) ub(9) extent(9) startIdx(1)

Traits: AttrSizedOperandSegments

Interfaces: NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes: 

AttributeMLIR TypeDescription
strideInBytes::mlir::BoolAttrbool attribute

Operands: 

OperandDescription
lowerboundinteger or index
upperboundinteger or index
extentinteger or index
strideinteger or index
startIdxinteger or index

Results: 

ResultDescription
resultType for representing acc data clause bounds information

acc.cache (acc::CacheOp) 

Represents the cache directive that is associated with a loop.

Syntax:

operation ::= `acc.cache` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.copyin (acc::CopyinOp) 

Represents copyin semantics for acc data clauses like acc copyin and acc copy.

Syntax:

operation ::= `acc.copyin` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.copyout (acc::CopyoutOp) 

Represents acc copyout semantics - reverse of copyin.

Syntax:

operation ::= `acc.copyout` `accPtr` `(` $accPtr `:` type($accPtr) `)`
              oilist(
              `bounds` `(` $bounds `)`
              | `to` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              ) attr-dict
  • varPtr: The address of variable to copy back to. This only applies to acc.copyout
  • accPtr: The acc address of variable. This is the link from the data-entry operation used.
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
accPtrPointerLikeType instance
varPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

acc.create (acc::CreateOp) 

Represents create semantics for acc data clauses like acc create and acc copyout.

Syntax:

operation ::= `acc.create` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.data (acc::DataOp) 

Data construct

Syntax:

operation ::= `acc.data` oilist(
              `if` `(` $ifCond `)`
              | `async` `(` $async `:` type($async) `)`
              | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              )
              $region attr-dict-with-keyword

The “acc.data” operation represents a data construct. It defines vars to be allocated in the current device memory for the duration of the region, whether data should be copied from local memory to the current device memory upon region entry , and copied from device memory to local memory upon region exit.

Example:

acc.data present(%a: memref<10x10xf32>, %b: memref<10x10xf32>,
    %c: memref<10xf32>, %d: memref<10xf32>) {
  // data region
}

Traits: AttrSizedOperandSegments, RecursiveMemoryEffects

Attributes: 

AttributeMLIR TypeDescription
asyncAttr::mlir::UnitAttrunit attribute
waitAttr::mlir::UnitAttrunit attribute
defaultAttr::mlir::acc::ClauseDefaultValueAttrDefaultValue Clause

Operands: 

OperandDescription
ifCond1-bit signless integer
asyncinteger or index
waitDevnuminteger or index
waitOperandsinteger or index
dataClauseOperandsPointerLikeType instance

acc.declare (acc::DeclareOp) 

Declare implicit region

Syntax:

operation ::= `acc.declare` `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              $region attr-dict-with-keyword

The “acc.declare” operation represents an implicit declare region in function (and subroutine in Fortran).

Example:

%pa = acc.present varPtr(%a : memref<10x10xf32>) -> memref<10x10xf32>
acc.declare dataOperands(%pa: memref<10x10xf32>) {
  // implicit region
}

Traits: RecursiveMemoryEffects

Operands: 

OperandDescription
dataClauseOperandsPointerLikeType instance

acc.declare_device_resident (acc::DeclareDeviceResidentOp) 

Represents acc declare device_resident semantics.

Syntax:

operation ::= `acc.declare_device_resident` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.declare_enter (acc::DeclareEnterOp) 

Declare directive - entry to implicit data region

Syntax:

operation ::= `acc.declare_enter` oilist(
              `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              attr-dict-with-keyword

The “acc.declare_enter” operation represents the OpenACC declare directive and captures the entry semantics to the implicit data region. This operation is modeled similarly to “acc.enter_data”.

Example showing acc declare create(a):

%0 = acc.create varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
acc.declare_enter dataOperands(%0 : !llvm.ptr<f32>)

Operands: 

OperandDescription
dataClauseOperandsPointerLikeType instance

acc.declare_exit (acc::DeclareExitOp) 

Declare directive - exit from implicit data region

Syntax:

operation ::= `acc.declare_exit` oilist(
              `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              attr-dict-with-keyword

The “acc.declare_exit” operation represents the OpenACC declare directive and captures the exit semantics from the implicit data region. This operation is modeled similarly to “acc.exit_data”.

Example showing acc declare device_resident(a):

%0 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause declare_device_resident>}
acc.declare_exit dataOperands(%0 : !llvm.ptr<f32>)
acc.delete accPtr(%0 : !llvm.ptr<f32>) {dataClause = #acc<data_clause declare_device_resident>}

Operands: 

OperandDescription
dataClauseOperandsPointerLikeType instance

Represents acc declare link semantics.

Syntax:

operation ::= `acc.declare_link` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.delete (acc::DeleteOp) 

Represents acc delete semantics - reverse of create.

Syntax:

operation ::= `acc.delete` `accPtr` `(` $accPtr `:` type($accPtr) `)`
              oilist(
              `bounds` `(` $bounds `)`
              | `to` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              ) attr-dict
  • varPtr: The address of variable to copy back to. This only applies to acc.copyout
  • accPtr: The acc address of variable. This is the link from the data-entry operation used.
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
accPtrPointerLikeType instance
varPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

acc.detach (acc::DetachOp) 

Represents acc detach semantics - reverse of attach.

Syntax:

operation ::= `acc.detach` `accPtr` `(` $accPtr `:` type($accPtr) `)`
              oilist(
              `bounds` `(` $bounds `)`
              | `to` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              ) attr-dict
  • varPtr: The address of variable to copy back to. This only applies to acc.copyout
  • accPtr: The acc address of variable. This is the link from the data-entry operation used.
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
accPtrPointerLikeType instance
varPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

acc.deviceptr (acc::DevicePtrOp) 

Specifies that the variable pointer is a device pointer.

Syntax:

operation ::= `acc.deviceptr` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.enter_data (acc::EnterDataOp) 

Enter data operation

Syntax:

operation ::= `acc.enter_data` oilist(
              `if` `(` $ifCond `)`
              | `async` `(` $asyncOperand `:` type($asyncOperand) `)`
              | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              attr-dict-with-keyword

The “acc.enter_data” operation represents the OpenACC enter data directive.

Example:

acc.enter_data create(%d1 : memref<10xf32>) attributes {async}

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
async::mlir::UnitAttrunit attribute
wait::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
ifCond1-bit signless integer
asyncOperandinteger or index
waitDevnuminteger or index
waitOperandsinteger or index
dataClauseOperandsPointerLikeType instance

acc.exit_data (acc::ExitDataOp) 

Exit data operation

Syntax:

operation ::= `acc.exit_data` oilist(
              `if` `(` $ifCond `)`
              | `async` `(` $asyncOperand `:` type($asyncOperand) `)`
              | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              attr-dict-with-keyword

The “acc.exit_data” operation represents the OpenACC exit data directive.

Example:

acc.exit_data delete(%d1 : memref<10xf32>) attributes {async}

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
async::mlir::UnitAttrunit attribute
wait::mlir::UnitAttrunit attribute
finalize::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
ifCond1-bit signless integer
asyncOperandinteger or index
waitDevnuminteger or index
waitOperandsinteger or index
dataClauseOperandsPointerLikeType instance

acc.firstprivate (acc::FirstprivateOp) 

Represents firstprivate semantic for the acc firstprivate clause.

Syntax:

operation ::= `acc.firstprivate` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.firstprivate.recipe (acc::FirstprivateRecipeOp) 

Privatization recipe

Syntax:

operation ::= `acc.firstprivate.recipe` $sym_name `:` $type attr-dict-with-keyword `init` $initRegion
              `copy` $copyRegion
              (`destroy` $destroyRegion^)?

Declares an OpenACC privatization recipe with copy of the initial value. The operation requires two mandatory regions and one optional.

  1. The initializer region specifies how to allocate and initialize a new private value. For example in Fortran, a derived-type might have a default initialization. The region has an argument that contains the value that need to be privatized. This is useful if the type is not known at compile time and the private value is needed to create its copy.
  2. The copy region specifies how to copy the initial value to the newly created private value. It takes the initial value and the privatized value as arguments.
  3. The destroy region specifies how to destruct the value when it reaches its end of life. It takes the privatized value as argument. It is optional.

A single privatization recipe can be used for multiple operand if they have the same type and do not require a specific default initialization.

Example:

acc.firstprivate.recipe @privatization_f32 : f32 init {
^bb0(%0: f32):
  // init region contains a sequence of operations to create and
  // initialize the copy if needed. It yields the create copy.
} copy {
^bb0(%0: f32, %1: !llvm.ptr<f32>):
  // copy region contains a sequence of operations to copy the initial value
  // of the firstprivate value to the newly created value.
} destroy {
^bb0(%0: f32)
  // destroy region contains a sequences of operations to destruct the
  // created copy.
}

// The privatization symbol is then used in the corresponding operation.
acc.parallel firstprivate(@privatization_f32 -> %a : f32) {
}

Traits: IsolatedFromAbove

Interfaces: Symbol

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute
type::mlir::TypeAttrany type attribute

acc.getdeviceptr (acc::GetDevicePtrOp) 

Gets device address if variable exists on device.

Syntax:

operation ::= `acc.getdeviceptr` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

This operation is used to get the accPtr for a variable. This is often used in conjunction with data exit operations when the data entry operation is not visible. This operation can have a dataClause argument that is any of the valid mlir::acc::DataClause entries. \

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.global_ctor (acc::GlobalConstructorOp) 

Used to hold construction operations associated with globals such as declare

Syntax:

operation ::= `acc.global_ctor` $sym_name $region attr-dict-with-keyword

The “acc.global_ctor” operation is used to capture OpenACC actions to apply on globals (such as acc declare) at the entry to the implicit data region. This operation is isolated and intended to be used in a module.

Example showing declare create of global:

llvm.mlir.global external @globalvar() : i32 {
  %0 = llvm.mlir.constant(0 : i32) : i32
  llvm.return %0 : i32
}
acc.global_ctor @acc_constructor {
  %0 = llvm.mlir.addressof @globalvar : !llvm.ptr<i32>
  %1 = acc.create varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32>
  acc.declare_enter dataOperands(%1 : !llvm.ptr<i32>)
}

Traits: IsolatedFromAbove

Interfaces: Symbol

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute

acc.global_dtor (acc::GlobalDestructorOp) 

Used to hold destruction operations associated with globals such as declare

Syntax:

operation ::= `acc.global_dtor` $sym_name $region attr-dict-with-keyword

The “acc.global_dtor” operation is used to capture OpenACC actions to apply on globals (such as acc declare) at the exit from the implicit data region. This operation is isolated and intended to be used in a module.

Example showing delete associated with declare create of global:

llvm.mlir.global external @globalvar() : i32 {
  %0 = llvm.mlir.constant(0 : i32) : i32
  llvm.return %0 : i32
}
acc.global_dtor @acc_destructor {
  %0 = llvm.mlir.addressof @globalvar : !llvm.ptr<i32>
  %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32> {dataClause = #acc<data_clause create>}
  acc.declare_exit dataOperands(%1 : !llvm.ptr<i32>)
  acc.delete accPtr(%1 : !llvm.ptr<i32>) {dataClause = #acc<data_clause create>}
}

Traits: IsolatedFromAbove

Interfaces: Symbol

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute

acc.host_data (acc::HostDataOp) 

Host_data construct

Syntax:

operation ::= `acc.host_data` oilist(
              `if` `(` $ifCond `)`
              | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              $region attr-dict-with-keyword

The “acc.host_data” operation represents the OpenACC host_data construct.

Example:

%0 = acc.use_device varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
acc.host_data dataOperands(%0 : !llvm.ptr<f32>) {

}

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
ifPresent::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
ifCond1-bit signless integer
dataClauseOperandsPointerLikeType instance

acc.init (acc::InitOp) 

Init operation

Syntax:

operation ::= `acc.init` oilist(
              `device_type` `(` $deviceTypeOperands `:` type($deviceTypeOperands) `)`
              | `device_num` `(` $deviceNumOperand `:` type($deviceNumOperand) `)`
              | `if` `(` $ifCond `)`
              ) attr-dict-with-keyword

The “acc.init” operation represents the OpenACC init executable directive.

Example:

acc.init
acc.init device_num(%dev1 : i32)

Traits: AttrSizedOperandSegments

Operands: 

OperandDescription
deviceTypeOperandsinteger
deviceNumOperandinteger or index
ifCond1-bit signless integer

acc.kernels (acc::KernelsOp) 

Kernels construct

Syntax:

operation ::= `acc.kernels` oilist(
              `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              | `async` `(` $async `:` type($async) `)`
              | `num_gangs` `(` $numGangs `:` type($numGangs) `)`
              | `num_workers` `(` $numWorkers `:` type($numWorkers) `)`
              | `vector_length` `(` $vectorLength `:` type($vectorLength) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `self` `(` $selfCond `)`
              | `if` `(` $ifCond `)`
              )
              $region attr-dict-with-keyword

The “acc.kernels” operation represents a kernels construct block. It has one region to be compiled into a sequence of kernels for execution on the current device.

Example:

acc.kernels num_gangs(%c10) num_workers(%c10)
    private(%c : memref<10xf32>) {
  // kernels region
}

Traits: AttrSizedOperandSegments, RecursiveMemoryEffects

Attributes: 

AttributeMLIR TypeDescription
asyncAttr::mlir::UnitAttrunit attribute
waitAttr::mlir::UnitAttrunit attribute
selfAttr::mlir::UnitAttrunit attribute
defaultAttr::mlir::acc::ClauseDefaultValueAttrDefaultValue Clause

Operands: 

OperandDescription
asyncinteger or index
waitOperandsinteger or index
numGangsinteger or index
numWorkersinteger or index
vectorLengthinteger or index
ifCond1-bit signless integer
selfCond1-bit signless integer
dataClauseOperandsPointerLikeType instance

acc.loop (acc::LoopOp) 

Loop construct

Syntax:

operation ::= `acc.loop` oilist(
              `gang` `` custom<GangClause>($gangNum, type($gangNum), $gangDim, type($gangDim), $gangStatic, type($gangStatic), $hasGang)
              | `worker` `` custom<WorkerClause>($workerNum, type($workerNum), $hasWorker)
              | `vector` `` custom<VectorClause>($vectorLength, type($vectorLength), $hasVector)
              | `private` `(` custom<SymOperandList>(
              $privateOperands, type($privateOperands), $privatizations)
              `)`
              | `tile` `(` $tileOperands `:` type($tileOperands) `)`
              | `reduction` `(` custom<SymOperandList>(
              $reductionOperands, type($reductionOperands), $reductionRecipes)
              `)`
              | `cache` `(` $cacheOperands `:` type($cacheOperands) `)`
              )
              $region
              ( `(` type($results)^ `)` )?
              attr-dict-with-keyword

The “acc.loop” operation represents the OpenACC loop construct.

Example:

acc.loop gang vector {
  scf.for %arg3 = %c0 to %c10 step %c1 {
    scf.for %arg4 = %c0 to %c10 step %c1 {
      scf.for %arg5 = %c0 to %c10 step %c1 {
        // ... body
      }
    }
  }
  acc.yield
} attributes { collapse = 3 }

Traits: AttrSizedOperandSegments, RecursiveMemoryEffects

Attributes: 

AttributeMLIR TypeDescription
collapse::mlir::IntegerAttr64-bit signless integer attribute
seq::mlir::UnitAttrunit attribute
independent::mlir::UnitAttrunit attribute
auto_::mlir::UnitAttrunit attribute
hasGang::mlir::UnitAttrunit attribute
hasWorker::mlir::UnitAttrunit attribute
hasVector::mlir::UnitAttrunit attribute
privatizations::mlir::ArrayAttrsymbol ref array attribute
reductionRecipes::mlir::ArrayAttrsymbol ref array attribute

Operands: 

OperandDescription
gangNuminteger or index
gangDiminteger or index
gangStaticinteger or index
workerNuminteger or index
vectorLengthinteger or index
tileOperandsinteger or index
privateOperandsPointerLikeType instance
reductionOperandsany type
cacheOperandsPointerLikeType instance

Results: 

ResultDescription
resultsany type

acc.nocreate (acc::NoCreateOp) 

Represents acc no_create semantics.

Syntax:

operation ::= `acc.nocreate` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.parallel (acc::ParallelOp) 

Parallel construct

Syntax:

operation ::= `acc.parallel` oilist(
              `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              | `async` `(` $async `:` type($async) `)`
              | `firstprivate` `(` custom<SymOperandList>($gangFirstPrivateOperands,
              type($gangFirstPrivateOperands), $firstprivatizations)
              `)`
              | `num_gangs` `(` $numGangs `:` type($numGangs) `)`
              | `num_workers` `(` $numWorkers `:` type($numWorkers) `)`
              | `private` `(` custom<SymOperandList>(
              $gangPrivateOperands, type($gangPrivateOperands), $privatizations)
              `)`
              | `vector_length` `(` $vectorLength `:` type($vectorLength) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `self` `(` $selfCond `)`
              | `if` `(` $ifCond `)`
              | `reduction` `(` custom<SymOperandList>(
              $reductionOperands, type($reductionOperands), $reductionRecipes)
              `)`
              )
              $region attr-dict-with-keyword

The “acc.parallel” operation represents a parallel construct block. It has one region to be executed in parallel on the current device.

Example:

acc.parallel num_gangs(%c10) num_workers(%c10)
    private(%c : memref<10xf32>) {
  // parallel region
}

Traits: AttrSizedOperandSegments, RecursiveMemoryEffects

Attributes: 

AttributeMLIR TypeDescription
asyncAttr::mlir::UnitAttrunit attribute
waitAttr::mlir::UnitAttrunit attribute
selfAttr::mlir::UnitAttrunit attribute
reductionRecipes::mlir::ArrayAttrsymbol ref array attribute
privatizations::mlir::ArrayAttrsymbol ref array attribute
firstprivatizations::mlir::ArrayAttrsymbol ref array attribute
defaultAttr::mlir::acc::ClauseDefaultValueAttrDefaultValue Clause

Operands: 

OperandDescription
asyncinteger or index
waitOperandsinteger or index
numGangsinteger or index
numWorkersinteger or index
vectorLengthinteger or index
ifCond1-bit signless integer
selfCond1-bit signless integer
reductionOperandsany type
gangPrivateOperandsPointerLikeType instance
gangFirstPrivateOperandsPointerLikeType instance
dataClauseOperandsPointerLikeType instance

acc.present (acc::PresentOp) 

Specifies that the variable is already present on device.

Syntax:

operation ::= `acc.present` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.private (acc::PrivateOp) 

Represents private semantics for acc private clause.

Syntax:

operation ::= `acc.private` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.private.recipe (acc::PrivateRecipeOp) 

Privatization recipe

Syntax:

operation ::= `acc.private.recipe` $sym_name `:` $type attr-dict-with-keyword `init` $initRegion
              (`destroy` $destroyRegion^)?

Declares an OpenACC privatization recipe. The operation requires one mandatory and one optional region.

  1. The initializer region specifies how to allocate and initialize a new private value. For example in Fortran, a derived-type might have a default initialization. The region has an argument that contains the value that need to be privatized. This is useful if the type is not known at compile time and the private value is needed to create its copy.
  2. The destroy region specifies how to destruct the value when it reaches its end of life. It takes the privatized value as argument.

A single privatization recipe can be used for multiple operand if they have the same type and do not require a specific default initialization.

Example:

acc.private.recipe @privatization_f32 : f32 init {
^bb0(%0: f32):
  // init region contains a sequence of operations to create and
  // initialize the copy if needed. It yields the create copy.
} destroy {
^bb0(%0: f32)
  // destroy region contains a sequences of operations to destruct the
  // created copy.
}

// The privatization symbol is then used in the corresponding operation.
acc.parallel private(@privatization_f32 -> %a : f32) {
}

Traits: IsolatedFromAbove

Interfaces: Symbol

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute
type::mlir::TypeAttrany type attribute

acc.reduction (acc::ReductionOp) 

Represents reduction semantics for acc reduction clause.

Syntax:

operation ::= `acc.reduction` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.reduction.recipe (acc::ReductionRecipeOp) 

Reduction recipe

Syntax:

operation ::= `acc.reduction.recipe` $sym_name `:` $type attr-dict-with-keyword
              `reduction_operator` $reductionOperator
              `init` $initRegion `combiner` $combinerRegion

Declares an OpenACC reduction recipe. The operation requires two mandatory regions.

  1. The initializer region specifies how to initialize the local reduction value. The region has a first argument that contains the value of the reduction accumulator at the start of the reduction. It is expected to acc.yield the new value. Extra arguments can be added to deal with dynamic arrays.
  2. The reduction region contains a sequences of operations to combine two values of the reduction type into one. It has at least two arguments and it is expected to acc.yield the combined value. Extra arguments can be added to deal with dynamic arrays.

Example:

acc.reduction.recipe @reduction_add_i64 : i64 reduction_operator<add> init {
^bb0(%0: i64):
  // init region contains a sequence of operations to initialize the local
  // reduction value as specified in 2.5.15
  %c0 = arith.constant 0 : i64
  acc.yield %c0 : i64
} combiner {
^bb0(%0: i64, %1: i64)
  // combiner region contains a sequence of operations to combine
  // two values into one.
  %2 = arith.addi %0, %1 : i64
  acc.yield %2 : i64
}

// The reduction symbol is then used in the corresponding operation.
acc.parallel reduction(@reduction_add_i64 -> %a : i64) {
}

The following table lists the valid operators and the initialization values according to OpenACC 3.3:

|————————————————|

C/C++Fortran
operatorinit value
+0
*1
maxleast
minlargest
&~0
^0
&&1
————————————————-

Traits: IsolatedFromAbove

Interfaces: Symbol

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute
type::mlir::TypeAttrany type attribute
reductionOperator::mlir::acc::ReductionOperatorAttrbuilt-in reduction operations supported by OpenACC

acc.routine (acc::RoutineOp) 

Acc routine operation

Syntax:

operation ::= `acc.routine` $sym_name `func` `(` $func_name `)`
              oilist (
              `bind` `(` $bind_name `)`
              | `gang` `` custom<RoutineGangClause>($gang, $gangDim)
              | `worker` $worker
              | `vector` $vector
              | `seq` $seq
              | `nohost` $nohost
              | `implicit` $implicit
              ) attr-dict-with-keyword

The acc.routine operation is used to capture the clauses of acc routine directive, including the associated function name. The associated function keeps track of its corresponding routine declaration through the RoutineInfoAttr.

Example:

func.func @acc_func(%a : i64) -> () attributes {acc.routine_info = #acc.routine_info<[@acc_func_rout1]>} {
  return
}
acc.routine @acc_func_rout1 func(@acc_func) gang

Traits: IsolatedFromAbove

Attributes: 

AttributeMLIR TypeDescription
sym_name::mlir::StringAttrstring attribute
func_name::mlir::StringAttrstring attribute
bind_name::mlir::StringAttrstring attribute
gang::mlir::UnitAttrunit attribute
worker::mlir::UnitAttrunit attribute
vector::mlir::UnitAttrunit attribute
seq::mlir::UnitAttrunit attribute
nohost::mlir::UnitAttrunit attribute
implicit::mlir::UnitAttrunit attribute
gangDim::mlir::IntegerAttrarbitrary integer attribute

acc.serial (acc::SerialOp) 

Serial construct

Syntax:

operation ::= `acc.serial` oilist(
              `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              | `async` `(` $async `:` type($async) `)`
              | `firstprivate` `(` custom<SymOperandList>($gangFirstPrivateOperands,
              type($gangFirstPrivateOperands), $firstprivatizations)
              `)`
              | `private` `(` custom<SymOperandList>(
              $gangPrivateOperands, type($gangPrivateOperands), $privatizations)
              `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `self` `(` $selfCond `)`
              | `if` `(` $ifCond `)`
              | `reduction` `(` custom<SymOperandList>(
              $reductionOperands, type($reductionOperands), $reductionRecipes)
              `)`
              )
              $region attr-dict-with-keyword

The “acc.serial” operation represents a serial construct block. It has one region to be executed in serial on the current device.

Example:

acc.serial private(%c : memref<10xf32>) {
  // serial region
}

Traits: AttrSizedOperandSegments, RecursiveMemoryEffects

Attributes: 

AttributeMLIR TypeDescription
asyncAttr::mlir::UnitAttrunit attribute
waitAttr::mlir::UnitAttrunit attribute
selfAttr::mlir::UnitAttrunit attribute
reductionRecipes::mlir::ArrayAttrsymbol ref array attribute
privatizations::mlir::ArrayAttrsymbol ref array attribute
firstprivatizations::mlir::ArrayAttrsymbol ref array attribute
defaultAttr::mlir::acc::ClauseDefaultValueAttrDefaultValue Clause

Operands: 

OperandDescription
asyncinteger or index
waitOperandsinteger or index
ifCond1-bit signless integer
selfCond1-bit signless integer
reductionOperandsany type
gangPrivateOperandsPointerLikeType instance
gangFirstPrivateOperandsPointerLikeType instance
dataClauseOperandsPointerLikeType instance

acc.set (acc::SetOp) 

Set operation

Syntax:

operation ::= `acc.set` oilist(
              `device_type` `(` $deviceType `:` type($deviceType) `)`
              | `default_async` `(` $defaultAsync `:` type($defaultAsync) `)`
              | `device_num` `(` $deviceNum `:` type($deviceNum) `)`
              | `if` `(` $ifCond `)`
              ) attr-dict-with-keyword

The “acc.set” operation represents the OpenACC set directive.

Example:

acc.set device_num(%dev1 : i32)

Traits: AttrSizedOperandSegments

Operands: 

OperandDescription
deviceTypeinteger or index
defaultAsyncinteger or index
deviceNuminteger or index
ifCond1-bit signless integer

acc.shutdown (acc::ShutdownOp) 

Shutdown operation

Syntax:

operation ::= `acc.shutdown` oilist(`device_type` `(` $deviceTypeOperands `:` type($deviceTypeOperands) `)`
              |`device_num` `(` $deviceNumOperand `:` type($deviceNumOperand) `)`
              |`if` `(` $ifCond `)`
              ) attr-dict-with-keyword

The “acc.shutdown” operation represents the OpenACC shutdown executable directive.

Example:

acc.shutdown
acc.shutdown device_num(%dev1 : i32)

Traits: AttrSizedOperandSegments

Operands: 

OperandDescription
deviceTypeOperandsinteger
deviceNumOperandinteger or index
ifCond1-bit signless integer

acc.terminator (acc::TerminatorOp) 

Generic terminator for OpenACC regions

Syntax:

operation ::= `acc.terminator` attr-dict

A terminator operation for regions that appear in the body of OpenACC operation. Generic OpenACC construct regions are not expected to return any value so the terminator takes no operands. The terminator op returns control to the enclosing op.

Traits: Terminator

acc.update (acc::UpdateOp) 

Update operation

Syntax:

operation ::= `acc.update` oilist(
              `if` `(` $ifCond `)`
              | `async` `(` $asyncOperand `:` type($asyncOperand) `)`
              | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)`
              | `device_type` `(` $deviceTypeOperands `:`
              type($deviceTypeOperands) `)`
              | `wait` `(` $waitOperands `:` type($waitOperands) `)`
              | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)`
              )
              attr-dict-with-keyword

The “acc.udpate” operation represents the OpenACC update executable directive. As host and self clauses are synonyms, any operands for host and self are add to $hostOperands.

Example:

acc.update device(%d1 : memref<10xf32>) attributes {async}

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
async::mlir::UnitAttrunit attribute
wait::mlir::UnitAttrunit attribute
ifPresent::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
ifCond1-bit signless integer
asyncOperandinteger or index
waitDevnuminteger or index
waitOperandsinteger or index
deviceTypeOperandsinteger or index
dataClauseOperandsPointerLikeType instance

acc.update_device (acc::UpdateDeviceOp) 

Represents acc update device semantics.

Syntax:

operation ::= `acc.update_device` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.update_host (acc::UpdateHostOp) 

Represents acc update host semantics.

Syntax:

operation ::= `acc.update_host` `accPtr` `(` $accPtr `:` type($accPtr) `)`
              oilist(
              `bounds` `(` $bounds `)`
              | `to` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              ) attr-dict
  • varPtr: The address of variable to copy back to. This only applies to acc.copyout
  • accPtr: The acc address of variable. This is the link from the data-entry operation used.
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
accPtrPointerLikeType instance
varPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

acc.use_device (acc::UseDeviceOp) 

Represents acc use_device semantics.

Syntax:

operation ::= `acc.use_device` `varPtr` `(` $varPtr `:` type($varPtr) `)`
              oilist(
              `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)`
              | `bounds` `(` $bounds `)`
              ) `->` type($accPtr) attr-dict

Description of arguments:

  • varPtr: The address of variable to copy.
  • varPtrPtr: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).
  • bounds: Used when copying just slice of array or array’s bounds are not encoded in type. They are in rank order where rank 0 is inner-most dimension.
  • dataClause: Keeps track of the data clause the user used. This is because the acc operations are decomposed. So a ‘copy’ clause is decomposed to both acc.copyin and acc.copyout operations, but both have dataClause that specifies acc_copy in this field.
  • structured: Flag to note whether this is associated with structured region (parallel, kernels, data) or unstructured (enter data, exit data). This is important due to spec specifically calling out structured and dynamic reference counters (2.6.7).
  • implicit: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.
  • name: Holds the name of variable as specified in user clause (including bounds).

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
dataClause::mlir::acc::DataClauseAttrdata clauses supported by OpenACC
structured::mlir::BoolAttrbool attribute
implicit::mlir::BoolAttrbool attribute
name::mlir::StringAttrstring attribute

Operands: 

OperandDescription
varPtrPointerLikeType instance
varPtrPtrPointerLikeType instance
boundsType for representing acc data clause bounds information

Results: 

ResultDescription
accPtrPointerLikeType instance

acc.wait (acc::WaitOp) 

Wait operation

Syntax:

operation ::= `acc.wait` ( `(` $waitOperands^ `:` type($waitOperands) `)` )?
              oilist(`async` `(` $asyncOperand `:` type($asyncOperand) `)`
              |`wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)`
              |`if` `(` $ifCond `)`
              ) attr-dict-with-keyword

The “acc.wait” operation represents the OpenACC wait executable directive.

Example:

acc.wait(%value1: index)
acc.wait() async(%async1: i32)

Traits: AttrSizedOperandSegments

Attributes: 

AttributeMLIR TypeDescription
async::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
waitOperandsinteger or index
asyncOperandinteger or index
waitDevnuminteger or index
ifCond1-bit signless integer

acc.yield (acc::YieldOp) 

Acc yield and termination operation

Syntax:

operation ::= `acc.yield` attr-dict ($operands^ `:` type($operands))?

acc.yield is a special terminator operation for block inside regions in various acc ops (including parallel, loop, atomic.update). It returns values to the immediately enclosing acc op.

Traits: HasParent<FirstprivateRecipeOp, LoopOp, ParallelOp, PrivateRecipeOp,ReductionRecipeOp, SerialOp, AtomicUpdateOp>, ReturnLike, Terminator

Interfaces: RegionBranchTerminatorOpInterface

Operands: 

OperandDescription
operandsany type

Attribute definition 

DeclareActionAttr 

Syntax:

#acc.declare_action<
  SymbolRefAttr,   # preAlloc
  SymbolRefAttr,   # postAlloc
  SymbolRefAttr,   # preDealloc
  SymbolRefAttr   # postDealloc
>

Parameters: 

ParameterC++ typeDescription
preAllocSymbolRefAttr
postAllocSymbolRefAttr
preDeallocSymbolRefAttr
postDeallocSymbolRefAttr

DeclareAttr 

Syntax:

#acc.declare<
  DataClauseAttr,   # dataClause
  bool   # implicit
>

Parameters: 

ParameterC++ typeDescription
dataClauseDataClauseAttr
implicitbool

ClauseDefaultValueAttr 

DefaultValue Clause

Syntax:

#acc.defaultvalue<
  ::mlir::acc::ClauseDefaultValue   # value
>

Parameters: 

ParameterC++ typeDescription
value::mlir::acc::ClauseDefaultValuean enum of type ClauseDefaultValue

DataClauseAttr 

data clauses supported by OpenACC

Syntax:

#acc.data_clause<
  ::mlir::acc::DataClause   # value
>

Parameters: 

ParameterC++ typeDescription
value::mlir::acc::DataClausean enum of type DataClause

ReductionOperatorAttr 

built-in reduction operations supported by OpenACC

Syntax:

#acc.reduction_operator<
  ::mlir::acc::ReductionOperator   # value
>

Parameters: 

ParameterC++ typeDescription
value::mlir::acc::ReductionOperatoran enum of type ReductionOperator

RoutineInfoAttr 

Keeps track of associated acc routine information

Syntax:

#acc.routine_info<
  ::llvm::ArrayRef<SymbolRefAttr>   # accRoutines
>

This attribute is used to create the association between a function and its acc.routine operation. A func.func uses this if its name was referenced in an acc routine directive.

Parameters: 

ParameterC++ typeDescription
accRoutines::llvm::ArrayRef<SymbolRefAttr>

Type definition 

DataBoundsType 

Type for representing acc data clause bounds information

Syntax: !acc.data_bounds_ty