mlir.dialects._acc_ops_gen

Attributes

_ods_ir

Classes

_Dialect
AtomicCaptureOp	This operation performs an atomic capture.
AtomicReadOp	This operation performs an atomic read.
AtomicUpdateOp	This operation performs an atomic update.
AtomicWriteOp	This operation performs an atomic write.
AttachOp	Description of arguments:
CacheOp	Description of arguments:
CopyinOp	Description of arguments:
CopyoutOp
CreateOp	Description of arguments:
DataBoundsOp	This operation is used to record bounds used in acc data clause in a
DataOp	The "acc.data" operation represents a data construct. It defines vars to
DeclareDeviceResidentOp	Description of arguments:
DeclareEnterOp	The "acc.declare_enter" operation represents the OpenACC declare directive
DeclareExitOp	The "acc.declare_exit" operation represents the OpenACC declare directive
DeclareLinkOp	Description of arguments:
DeclareOp	The "acc.declare" operation represents an implicit declare region in
DeleteOp
DetachOp
DevicePtrOp	Description of arguments:
EnterDataOp	The "acc.enter_data" operation represents the OpenACC enter data directive.
ExitDataOp	The "acc.exit_data" operation represents the OpenACC exit data directive.
FirstprivateMapInitialOp	Description of arguments:
FirstprivateOp	Description of arguments:
FirstprivateRecipeOp	Declares an OpenACC privatization recipe with copy of the initial value.
GetDevicePtrOp	This operation is used to get the accPtr for a variable. This is often
GetExtentOp	This operation extracts the extent value from an acc.bounds value.
GetLowerboundOp	This operation extracts the lowerbound value from an acc.bounds value.
GetStrideOp	This operation extracts the stride value from an acc.bounds value.
GetUpperboundOp	This operation extracts the upperbound value from an acc.bounds value.
GlobalConstructorOp	The "acc.global_ctor" operation is used to capture OpenACC actions to apply
GlobalDestructorOp	The "acc.global_dtor" operation is used to capture OpenACC actions to apply
HostDataOp	The "acc.host_data" operation represents the OpenACC host_data construct.
InitOp	The "acc.init" operation represents the OpenACC init executable
KernelEnvironmentOp	The acc.kernel_environment operation represents a decomposition of
KernelsOp	The "acc.kernels" operation represents a kernels construct block. It has
LoopOp	The acc.loop operation represents the OpenACC loop construct and when
NoCreateOp	Description of arguments:
ParallelOp	The "acc.parallel" operation represents a parallel construct block. It has
PresentOp	Description of arguments:
PrivateOp	Description of arguments:
PrivateRecipeOp	Declares an OpenACC privatization recipe. The operation requires one
ReductionOp	Description of arguments:
ReductionRecipeOp	Declares an OpenACC reduction recipe. The operation requires two
RoutineOp	The acc.routine operation is used to capture the clauses of acc
SerialOp	The "acc.serial" operation represents a serial construct block. It has
SetOp	The "acc.set" operation represents the OpenACC set directive.
ShutdownOp	The "acc.shutdown" operation represents the OpenACC shutdown executable
TerminatorOp	A terminator operation for regions that appear in the body of OpenACC
UpdateDeviceOp	Description of arguments:
UpdateHostOp
UpdateOp	The acc.update operation represents the OpenACC update executable
UseDeviceOp	Description of arguments:
WaitOp	The "acc.wait" operation represents the OpenACC wait executable
YieldOp	acc.yield is a special terminator operation for block inside regions in

Functions

atomic_capture(→ AtomicCaptureOp)
atomic_read(→ AtomicReadOp)
atomic_update(→ AtomicUpdateOp)
atomic_write(→ AtomicWriteOp)
attach(→ _ods_ir)
cache(→ _ods_ir)
copyin(→ _ods_ir)
copyout(→ CopyoutOp)
create_(→ _ods_ir)
bounds(→ _ods_ir)
data(→ DataOp)
declare_device_resident(→ _ods_ir)
declare_enter(→ _ods_ir)
declare_exit(→ DeclareExitOp)
declare_link(→ _ods_ir)
declare(→ DeclareOp)
delete(→ DeleteOp)
detach(→ DetachOp)
deviceptr(→ _ods_ir)
enter_data(→ EnterDataOp)
exit_data(→ ExitDataOp)
firstprivate_map(→ _ods_ir)
firstprivate(→ _ods_ir)
firstprivate_recipe(→ FirstprivateRecipeOp)
getdeviceptr(→ _ods_ir)
get_extent(→ _ods_ir)
get_lowerbound(→ _ods_ir)
get_stride(→ _ods_ir)
get_upperbound(→ _ods_ir)
global_ctor(→ GlobalConstructorOp)
global_dtor(→ GlobalDestructorOp)
host_data(→ HostDataOp)
init(→ InitOp)
kernel_environment(→ KernelEnvironmentOp)
kernels(→ KernelsOp)
loop(→ Union[_ods_ir, _ods_ir, LoopOp])
nocreate(→ _ods_ir)
parallel(→ ParallelOp)
present(→ _ods_ir)
private(→ _ods_ir)
private_recipe(→ PrivateRecipeOp)
reduction(→ _ods_ir)
reduction_recipe(→ ReductionRecipeOp)
routine(→ RoutineOp)
serial(→ SerialOp)
set(→ SetOp)
shutdown(→ ShutdownOp)
terminator(→ TerminatorOp)
update_device(→ _ods_ir)
update_host(→ UpdateHostOp)
update(→ UpdateOp)
use_device(→ _ods_ir)
wait(→ WaitOp)
yield_(→ YieldOp)

Module Contents

mlir.dialects._acc_ops_gen._ods_ir

class mlir.dialects._acc_ops_gen._Dialect(descriptor: object)

Bases: _ods_ir

DIALECT_NAMESPACE = 'acc'

class mlir.dialects._acc_ops_gen.AtomicCaptureOp(*, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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
}

OPERATION_NAME = 'acc.atomic.capture'

_ODS_REGIONS = (1, True)

ifCond() → _ods_ir | None

region() → _ods_ir

mlir.dialects._acc_ops_gen.atomic_capture(*, if_cond=None, loc=None, ip=None) → AtomicCaptureOp

class mlir.dialects._acc_ops_gen.AtomicReadOp(x, v, element_type, *, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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.

OPERATION_NAME = 'acc.atomic.read'

_ODS_REGIONS = (0, True)

x() → _ods_ir

v() → _ods_ir

ifCond() → _ods_ir | None

element_type() → _ods_ir

mlir.dialects._acc_ops_gen.atomic_read(x, v, element_type, *, if_cond=None, loc=None, ip=None) → AtomicReadOp

class mlir.dialects._acc_ops_gen.AtomicUpdateOp(x, *, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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.

OPERATION_NAME = 'acc.atomic.update'

_ODS_REGIONS = (1, True)

x() → _ods_ir

ifCond() → _ods_ir | None

region() → _ods_ir

mlir.dialects._acc_ops_gen.atomic_update(x, *, if_cond=None, loc=None, ip=None) → AtomicUpdateOp

class mlir.dialects._acc_ops_gen.AtomicWriteOp(x, expr, *, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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).

OPERATION_NAME = 'acc.atomic.write'

_ODS_REGIONS = (0, True)

x() → _ods_ir

expr() → _ods_ir

ifCond() → _ods_ir | None

mlir.dialects._acc_ops_gen.atomic_write(x, expr, *, if_cond=None, loc=None, ip=None) → AtomicWriteOp

class mlir.dialects._acc_ops_gen.AttachOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.attach'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.attach(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.CacheOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.cache'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.cache(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.CopyinOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.copyin'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.copyin(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.CopyoutOp(accVar, var, varType, bounds, asyncOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

• varPtr: The address of variable to copy back to.

• accVar: The acc 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, always, etc) * name: Holds the name of variable as specified in user clause (including bounds). The async values attached to the data exit operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.copyout'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

accVar() → _ods_ir

var() → _ods_ir

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

mlir.dialects._acc_ops_gen.copyout(acc_var, var, var_type, bounds, async_operands, *, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → CopyoutOp

class mlir.dialects._acc_ops_gen.CreateOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.create'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.create_(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.DataBoundsOp(result, *, lowerbound=None, upperbound=None, extent=None, stride=None, strideInBytes=None, startIdx=None, loc=None, ip=None)

Bases: _ods_ir

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.

The stride represents the distance between consecutive elements. For multi-dimensional arrays, the stride for each outer dimension must account for the complete size of all inner dimensions.

The strideInBytes flag indicates that the stride is specified in bytes rather than the number of elements.

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) stride(1)

Fortran:

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

=>

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

OPERATION_NAME = 'acc.bounds'

_ODS_OPERAND_SEGMENTS = [0, 0, 0, 0, 0]

_ODS_REGIONS = (0, True)

lowerbound() → _ods_ir | None

upperbound() → _ods_ir | None

extent() → _ods_ir | None

stride() → _ods_ir | None

startIdx() → _ods_ir | None

strideInBytes() → _ods_ir

result() → _ods_ir

Shortcut to get an op result if it has only one (throws an error otherwise).

mlir.dialects._acc_ops_gen.bounds(result, *, lowerbound=None, upperbound=None, extent=None, stride=None, stride_in_bytes=None, start_idx=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.DataOp(asyncOperands, waitOperands, dataClauseOperands, *, ifCond=None, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, defaultAttr=None, loc=None, ip=None)

Bases: _ods_ir

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
}

async and wait operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.data'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

ifCond() → _ods_ir | None

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

dataClauseOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

defaultAttr() → _ods_ir | None

region() → _ods_ir

mlir.dialects._acc_ops_gen.data(async_operands, wait_operands, data_clause_operands, *, if_cond=None, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, default_attr=None, loc=None, ip=None) → DataOp

class mlir.dialects._acc_ops_gen.DeclareDeviceResidentOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.declare_device_resident'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.declare_device_resident(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.DeclareEnterOp(token, dataClauseOperands, *, loc=None, ip=None)

Bases: _ods_ir

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) -> !llvm.ptr
acc.declare_enter dataOperands(%0 : !llvm.ptr)

OPERATION_NAME = 'acc.declare_enter'

_ODS_REGIONS = (0, True)

dataClauseOperands() → _ods_ir

token() → _ods_ir

mlir.dialects._acc_ops_gen.declare_enter(token, data_clause_operands, *, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.DeclareExitOp(dataClauseOperands, *, token=None, loc=None, ip=None)

Bases: _ods_ir

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) -> !llvm.ptr {dataClause = #acc<data_clause declare_device_resident>}
acc.declare_exit dataOperands(%0 : !llvm.ptr)
acc.delete accPtr(%0 : !llvm.ptr) {dataClause = #acc<data_clause declare_device_resident>}

OPERATION_NAME = 'acc.declare_exit'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

token() → _ods_ir | None

dataClauseOperands() → _ods_ir

mlir.dialects._acc_ops_gen.declare_exit(data_clause_operands, *, token=None, loc=None, ip=None) → DeclareExitOp

class mlir.dialects._acc_ops_gen.DeclareLinkOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.declare_link'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.declare_link(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.DeclareOp(dataClauseOperands, *, loc=None, ip=None)

Bases: _ods_ir

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
}

OPERATION_NAME = 'acc.declare'

_ODS_REGIONS = (1, True)

dataClauseOperands() → _ods_ir

region() → _ods_ir

mlir.dialects._acc_ops_gen.declare(data_clause_operands, *, loc=None, ip=None) → DeclareOp

class mlir.dialects._acc_ops_gen.DeleteOp(accVar, bounds, asyncOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

• accVar: The acc 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, always, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data exit operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.delete'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

accVar() → _ods_ir

bounds() → _ods_ir

asyncOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

mlir.dialects._acc_ops_gen.delete(acc_var, bounds, async_operands, *, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → DeleteOp

class mlir.dialects._acc_ops_gen.DetachOp(accVar, bounds, asyncOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

• accVar: The acc 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, always, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data exit operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.detach'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

accVar() → _ods_ir

bounds() → _ods_ir

asyncOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

mlir.dialects._acc_ops_gen.detach(acc_var, bounds, async_operands, *, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → DetachOp

class mlir.dialects._acc_ops_gen.DevicePtrOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.deviceptr'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.deviceptr(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.EnterDataOp(waitOperands, dataClauseOperands, *, ifCond=None, asyncOperand=None, async_=None, waitDevnum=None, wait=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

OPERATION_NAME = 'acc.enter_data'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

ifCond() → _ods_ir | None

asyncOperand() → _ods_ir | None

waitDevnum() → _ods_ir | None

waitOperands() → _ods_ir

dataClauseOperands() → _ods_ir

async_() → bool

wait() → bool

mlir.dialects._acc_ops_gen.enter_data(wait_operands, data_clause_operands, *, if_cond=None, async_operand=None, async_=None, wait_devnum=None, wait=None, loc=None, ip=None) → EnterDataOp

class mlir.dialects._acc_ops_gen.ExitDataOp(waitOperands, dataClauseOperands, *, ifCond=None, asyncOperand=None, async_=None, waitDevnum=None, wait=None, finalize=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

OPERATION_NAME = 'acc.exit_data'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

ifCond() → _ods_ir | None

asyncOperand() → _ods_ir | None

waitDevnum() → _ods_ir | None

waitOperands() → _ods_ir

dataClauseOperands() → _ods_ir

async_() → bool

wait() → bool

finalize() → bool

mlir.dialects._acc_ops_gen.exit_data(wait_operands, data_clause_operands, *, if_cond=None, async_operand=None, async_=None, wait_devnum=None, wait=None, finalize=None, loc=None, ip=None) → ExitDataOp

class mlir.dialects._acc_ops_gen.FirstprivateMapInitialOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.firstprivate_map'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.firstprivate_map(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.FirstprivateOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.firstprivate'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.firstprivate(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.FirstprivateRecipeOp(sym_name, type_, *, loc=None, ip=None)

Bases: _ods_ir

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

#. 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 original value that needs to be privatized, followed by bounds arguments (if any) in order from innermost to outermost dimension. The region must yield the privatized copy. #. The copy region specifies how to copy the initial value to the newly created private value. It takes the original value, the privatized value, followed by bounds arguments (if any) in the same order. #. The destroy region specifies how to destruct the value when it reaches its end of life. It takes the original value, the privatized value, and bounds arguments (if any) in the same order. 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 @firstprivate_memref : memref<10x20xf32> init {
^bb0(%original: memref<10x20xf32>):
  // init region contains a sequence of operations to create and
  // initialize the copy. It yields the privatized copy.
  %alloca = memref.alloca() : memref<10x20xf32>
  acc.yield %alloca : memref<10x20xf32>
} copy {
^bb0(%original: memref<10x20xf32>, %privatized: memref<10x20xf32>):
  // copy region contains a sequence of operations to copy the initial value
  // of the firstprivate value to the newly created value.
  memref.copy %original, %privatized : memref<10x20xf32> to memref<10x20xf32>
  acc.terminator
} destroy {
^bb0(%original: memref<10x20xf32>, %privatized: memref<10x20xf32>):
  // destroy region is empty since alloca is automatically cleaned up
  acc.terminator
}

// Example with bounds for array slicing:
acc.firstprivate.recipe @firstprivate_slice : memref<10x20xf32> init {
^bb0(%original: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Extract bounds and create appropriately sized allocation
  %extent_inner = acc.get_extent %bounds_inner : (!acc.data_bounds_ty) -> index
  %extent_outer = acc.get_extent %bounds_outer : (!acc.data_bounds_ty) -> index
  %slice_alloc = memref.alloca(%extent_outer, %extent_inner) : memref<?x?xf32>
  // ... base pointer adjustment logic ...
  acc.yield %result : memref<10x20xf32>
} copy {
^bb0(%original: memref<10x20xf32>, %privatized: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Copy the slice portion from original to privatized
  %lb_inner = acc.get_lowerbound %bounds_inner : (!acc.data_bounds_ty) -> index
  %lb_outer = acc.get_lowerbound %bounds_outer : (!acc.data_bounds_ty) -> index
  %extent_inner = acc.get_extent %bounds_inner : (!acc.data_bounds_ty) -> index
  %extent_outer = acc.get_extent %bounds_outer : (!acc.data_bounds_ty) -> index
  %subview = memref.subview %original[%lb_outer, %lb_inner][%extent_outer, %extent_inner][1, 1]
    : memref<10x20xf32> to memref<?x?xf32, strided<[20, 1], offset: ?>>
  // Copy subview to privatized...
  acc.terminator
}

// The privatization symbol is then used in the corresponding operation.
acc.parallel firstprivate(@firstprivate_memref -> %a : memref<10x20xf32>) {
}

OPERATION_NAME = 'acc.firstprivate.recipe'

_ODS_REGIONS = (3, True)

sym_name() → _ods_ir

type_() → _ods_ir

initRegion() → _ods_ir

copyRegion() → _ods_ir

destroyRegion() → _ods_ir

mlir.dialects._acc_ops_gen.firstprivate_recipe(sym_name, type_, *, loc=None, ip=None) → FirstprivateRecipeOp

class mlir.dialects._acc_ops_gen.GetDevicePtrOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

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:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.getdeviceptr'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.getdeviceptr(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.GetExtentOp(result, bounds, *, loc=None, ip=None)

Bases: _ods_ir

This operation extracts the extent value from an acc.bounds value. If the data bounds does not have an extent specified, it is computed from the upperbound.

Example:

%extent = acc.get_extent %bounds : (!acc.data_bounds_ty) -> index

OPERATION_NAME = 'acc.get_extent'

_ODS_REGIONS = (0, True)

bounds() → _ods_ir

result() → _ods_ir

Shortcut to get an op result if it has only one (throws an error otherwise).

mlir.dialects._acc_ops_gen.get_extent(result, bounds, *, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.GetLowerboundOp(result, bounds, *, loc=None, ip=None)

Bases: _ods_ir

This operation extracts the lowerbound value from an acc.bounds value. If the data bounds does not have a lowerbound specified, it means it is zero.

Example:

%lb = acc.get_lowerbound %bounds : (!acc.data_bounds_ty) -> index

OPERATION_NAME = 'acc.get_lowerbound'

_ODS_REGIONS = (0, True)

bounds() → _ods_ir

result() → _ods_ir

Shortcut to get an op result if it has only one (throws an error otherwise).

mlir.dialects._acc_ops_gen.get_lowerbound(result, bounds, *, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.GetStrideOp(result, bounds, *, loc=None, ip=None)

Bases: _ods_ir

This operation extracts the stride value from an acc.bounds value. If the data bounds does not have a stride specified, it defaults to 1.

Example:

%stride = acc.get_stride %bounds : (!acc.data_bounds_ty) -> index

OPERATION_NAME = 'acc.get_stride'

_ODS_REGIONS = (0, True)

bounds() → _ods_ir

result() → _ods_ir

Shortcut to get an op result if it has only one (throws an error otherwise).

mlir.dialects._acc_ops_gen.get_stride(result, bounds, *, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.GetUpperboundOp(result, bounds, *, loc=None, ip=None)

Bases: _ods_ir

This operation extracts the upperbound value from an acc.bounds value. If the data bounds does not have an upperbound specified, this operation uses the extent to compute it.

Example:

%ub = acc.get_upperbound %bounds : (!acc.data_bounds_ty) -> index

OPERATION_NAME = 'acc.get_upperbound'

_ODS_REGIONS = (0, True)

bounds() → _ods_ir

result() → _ods_ir

Shortcut to get an op result if it has only one (throws an error otherwise).

mlir.dialects._acc_ops_gen.get_upperbound(result, bounds, *, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.GlobalConstructorOp(sym_name, *, loc=None, ip=None)

Bases: _ods_ir

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
  %1 = acc.create varPtr(%0 : !llvm.ptr) -> !llvm.ptr
  acc.declare_enter dataOperands(%1 : !llvm.ptr)
}

OPERATION_NAME = 'acc.global_ctor'

_ODS_REGIONS = (1, True)

sym_name() → _ods_ir

region() → _ods_ir

mlir.dialects._acc_ops_gen.global_ctor(sym_name, *, loc=None, ip=None) → GlobalConstructorOp

class mlir.dialects._acc_ops_gen.GlobalDestructorOp(sym_name, *, loc=None, ip=None)

Bases: _ods_ir

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
  %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause create>}
  acc.declare_exit dataOperands(%1 : !llvm.ptr)
  acc.delete accPtr(%1 : !llvm.ptr) {dataClause = #acc<data_clause create>}
}

OPERATION_NAME = 'acc.global_dtor'

_ODS_REGIONS = (1, True)

sym_name() → _ods_ir

region() → _ods_ir

mlir.dialects._acc_ops_gen.global_dtor(sym_name, *, loc=None, ip=None) → GlobalDestructorOp

class mlir.dialects._acc_ops_gen.HostDataOp(dataClauseOperands, *, ifCond=None, ifPresent=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

}

OPERATION_NAME = 'acc.host_data'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

ifCond() → _ods_ir | None

dataClauseOperands() → _ods_ir

ifPresent() → bool

region() → _ods_ir

mlir.dialects._acc_ops_gen.host_data(data_clause_operands, *, if_cond=None, if_present=None, loc=None, ip=None) → HostDataOp

class mlir.dialects._acc_ops_gen.InitOp(*, device_types=None, deviceNum=None, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

OPERATION_NAME = 'acc.init'

_ODS_OPERAND_SEGMENTS = [0, 0]

_ODS_REGIONS = (0, True)

deviceNum() → _ods_ir | None

ifCond() → _ods_ir | None

device_types() → _ods_ir | None

mlir.dialects._acc_ops_gen.init(*, device_types=None, device_num=None, if_cond=None, loc=None, ip=None) → InitOp

class mlir.dialects._acc_ops_gen.KernelEnvironmentOp(dataClauseOperands, asyncOperands, waitOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, loc=None, ip=None)

Bases: _ods_ir

The acc.kernel_environment operation represents a decomposition of any OpenACC compute construct (acc.kernels, acc.parallel, or acc.serial) that captures data mapping and asynchronous behavior:

• data clause operands

• async clause operands

• wait clause operands

This allows kernel execution parallelism and privatization to be handled separately, facilitating eventual lowering to GPU dialect where kernel launching and compute offloading are handled separately.

OPERATION_NAME = 'acc.kernel_environment'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

dataClauseOperands() → _ods_ir

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

region() → _ods_ir

mlir.dialects._acc_ops_gen.kernel_environment(data_clause_operands, async_operands, wait_operands, *, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, loc=None, ip=None) → KernelEnvironmentOp

class mlir.dialects._acc_ops_gen.KernelsOp(asyncOperands, waitOperands, numGangs, numWorkers, vectorLength, dataClauseOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, numGangsSegments=None, numGangsDeviceType=None, numWorkersDeviceType=None, vectorLengthDeviceType=None, ifCond=None, selfCond=None, selfAttr=None, defaultAttr=None, combined=None, loc=None, ip=None)

Bases: _ods_ir

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
}

collapse, gang, worker, vector, seq, independent, auto and tile operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.kernels'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

numGangs() → _ods_ir

numWorkers() → _ods_ir

vectorLength() → _ods_ir

ifCond() → _ods_ir | None

selfCond() → _ods_ir | None

dataClauseOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

numGangsSegments() → _ods_ir | None

numGangsDeviceType() → _ods_ir | None

numWorkersDeviceType() → _ods_ir | None

vectorLengthDeviceType() → _ods_ir | None

selfAttr() → bool

defaultAttr() → _ods_ir | None

combined() → bool

region() → _ods_ir

mlir.dialects._acc_ops_gen.kernels(async_operands, wait_operands, num_gangs, num_workers, vector_length, data_clause_operands, *, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, num_gangs_segments=None, num_gangs_device_type=None, num_workers_device_type=None, vector_length_device_type=None, if_cond=None, self_cond=None, self_attr=None, default_attr=None, combined=None, loc=None, ip=None) → KernelsOp

class mlir.dialects._acc_ops_gen.LoopOp(results_, lowerbound, upperbound, step, gangOperands, workerNumOperands, vectorOperands, tileOperands, cacheOperands, privateOperands, firstprivateOperands, reductionOperands, *, inclusiveUpperbound=None, collapse=None, collapseDeviceType=None, gangOperandsArgType=None, gangOperandsSegments=None, gangOperandsDeviceType=None, workerNumOperandsDeviceType=None, vectorOperandsDeviceType=None, seq=None, independent=None, auto_=None, gang=None, worker=None, vector=None, tileOperandsSegments=None, tileOperandsDeviceType=None, privatizationRecipes=None, firstprivatizationRecipes=None, reductionRecipes=None, combined=None, unstructured=None, loc=None, ip=None)

Bases: _ods_ir

The acc.loop operation represents the OpenACC loop construct and when bounds are included, the associated source language loop iterators. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound. If the inclusive attribute is set then the upper bound is included.

In cases where the OpenACC loop directive needs to capture multiple source language loops, such as in the case of collapse or tile, the multiple induction arguments are used to capture each case. Having such a representation makes sure no intermediate transformation such as Loop Invariant Code Motion breaks the property requested by the clause on the loop constructs.

Each acc.loop holds private and reduction operands which are the ssa values from the corresponding acc.private or acc.reduction operations. Additionally, firstprivate operands are supported to represent cases where privatization is needed with initialization from an original value. While the OpenACC specification does not explicitly support firstprivate on loop constructs, this extension enables representing privatization scenarios that arise from an optimization and codegen pipeline operating on acc dialect.

The operation supports capturing information that it comes combined constructs (e.g., parallel loop, kernels loop, serial loop) through the combined attribute despite requiring the acc.loop to be decomposed from the compute operation representing compute construct.

Example:

acc.loop gang() vector() (%arg3 : index, %arg4 : index, %arg5 : index) =
    (%c0, %c0, %c0 : index, index, index) to
    (%c10, %c10, %c10 : index, index, index) step
    (%c1, %c1, %c1 : index, index, index) {
  // Loop body
  acc.yield
} attributes { collapse = [3] }

collapse, gang, worker, vector, seq, independent, auto, cache, and tile operands are supported with device_type information. These clauses should only be accessed through the provided device-type-aware getter methods. When modifying these operands, the corresponding device_type attributes must be updated to maintain consistency between operands and their target device types.

The unstructured attribute indicates that the loops inside the OpenACC construct contain early exits and cannot be lowered to structured MLIR operations. When this flag is set, the acc.loop should have no induction variables and the loop must be implemented via explicit control flow inside its body.

OPERATION_NAME = 'acc.loop'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

lowerbound() → _ods_ir

upperbound() → _ods_ir

step() → _ods_ir

gangOperands() → _ods_ir

workerNumOperands() → _ods_ir

vectorOperands() → _ods_ir

tileOperands() → _ods_ir

cacheOperands() → _ods_ir

privateOperands() → _ods_ir

firstprivateOperands() → _ods_ir

reductionOperands() → _ods_ir

inclusiveUpperbound() → _ods_ir | None

collapse() → _ods_ir | None

collapseDeviceType() → _ods_ir | None

gangOperandsArgType() → _ods_ir | None

gangOperandsSegments() → _ods_ir | None

gangOperandsDeviceType() → _ods_ir | None

workerNumOperandsDeviceType() → _ods_ir | None

vectorOperandsDeviceType() → _ods_ir | None

seq() → _ods_ir | None

independent() → _ods_ir | None

auto_() → _ods_ir | None

gang() → _ods_ir | None

worker() → _ods_ir | None

vector() → _ods_ir | None

tileOperandsSegments() → _ods_ir | None

tileOperandsDeviceType() → _ods_ir | None

privatizationRecipes() → _ods_ir | None

firstprivatizationRecipes() → _ods_ir | None

reductionRecipes() → _ods_ir | None

combined() → _ods_ir | None

unstructured() → bool

results_() → _ods_ir

region() → _ods_ir

mlir.dialects._acc_ops_gen.loop(results_, lowerbound, upperbound, step, gang_operands, worker_num_operands, vector_operands, tile_operands, cache_operands, private_operands, firstprivate_operands, reduction_operands, *, inclusive_upperbound=None, collapse=None, collapse_device_type=None, gang_operands_arg_type=None, gang_operands_segments=None, gang_operands_device_type=None, worker_num_operands_device_type=None, vector_operands_device_type=None, seq=None, independent=None, auto_=None, gang=None, worker=None, vector=None, tile_operands_segments=None, tile_operands_device_type=None, privatization_recipes=None, firstprivatization_recipes=None, reduction_recipes=None, combined=None, unstructured=None, loc=None, ip=None) → _ods_ir | _ods_ir | LoopOp

class mlir.dialects._acc_ops_gen.NoCreateOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.nocreate'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.nocreate(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.ParallelOp(asyncOperands, waitOperands, numGangs, numWorkers, vectorLength, reductionOperands, privateOperands, firstprivateOperands, dataClauseOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, numGangsSegments=None, numGangsDeviceType=None, numWorkersDeviceType=None, vectorLengthDeviceType=None, ifCond=None, selfCond=None, selfAttr=None, reductionRecipes=None, privatizationRecipes=None, firstprivatizationRecipes=None, defaultAttr=None, combined=None, loc=None, ip=None)

Bases: _ods_ir

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
}

async, wait, num_gangs, num_workers and vector_length operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.parallel'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

numGangs() → _ods_ir

numWorkers() → _ods_ir

vectorLength() → _ods_ir

ifCond() → _ods_ir | None

selfCond() → _ods_ir | None

reductionOperands() → _ods_ir

privateOperands() → _ods_ir

firstprivateOperands() → _ods_ir

dataClauseOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

numGangsSegments() → _ods_ir | None

numGangsDeviceType() → _ods_ir | None

numWorkersDeviceType() → _ods_ir | None

vectorLengthDeviceType() → _ods_ir | None

selfAttr() → bool

reductionRecipes() → _ods_ir | None

privatizationRecipes() → _ods_ir | None

firstprivatizationRecipes() → _ods_ir | None

defaultAttr() → _ods_ir | None

combined() → bool

region() → _ods_ir

mlir.dialects._acc_ops_gen.parallel(async_operands, wait_operands, num_gangs, num_workers, vector_length, reduction_operands, private_operands, firstprivate_operands, data_clause_operands, *, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, num_gangs_segments=None, num_gangs_device_type=None, num_workers_device_type=None, vector_length_device_type=None, if_cond=None, self_cond=None, self_attr=None, reduction_recipes=None, privatization_recipes=None, firstprivatization_recipes=None, default_attr=None, combined=None, loc=None, ip=None) → ParallelOp

class mlir.dialects._acc_ops_gen.PresentOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.present'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.present(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.PrivateOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.private'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.private(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.PrivateRecipeOp(sym_name, type_, *, loc=None, ip=None)

Bases: _ods_ir

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

#. 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 original value that needs to be privatized, followed by bounds arguments (if any) in order from innermost to outermost dimension. The region must yield the privatized copy. #. The destroy region specifies how to destruct the value when it reaches its end of life. It takes the original value, the privatized value, and bounds arguments (if any) in the same order as the init region.

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_memref : memref<10x20xf32> init {
^bb0(%original: memref<10x20xf32>):
  // init region contains a sequence of operations to create and
  // initialize the copy. It yields the privatized copy.
  %alloca = memref.alloca() : memref<10x20xf32>
  acc.yield %alloca : memref<10x20xf32>
} destroy {
^bb0(%original: memref<10x20xf32>, %privatized: memref<10x20xf32>):
  // destroy region is empty since alloca is automatically cleaned up
  acc.terminator
}

// Example with bounds for array slicing:
acc.private.recipe @privatization_slice : memref<10x20xf32> init {
^bb0(%original: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Extract bounds and create appropriately sized allocation
  %extent_inner = acc.get_extent %bounds_inner : (!acc.data_bounds_ty) -> index
  %extent_outer = acc.get_extent %bounds_outer : (!acc.data_bounds_ty) -> index
  %slice_alloc = memref.alloca(%extent_outer, %extent_inner) : memref<?x?xf32>
  // ... base pointer adjustment logic ...
  acc.yield %result : memref<10x20xf32>
} destroy {
^bb0(%original: memref<10x20xf32>, %privatized: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Cleanup is automatic for alloca-based allocations
  acc.terminator
}

// The privatization symbol is then used in the corresponding operation.
acc.parallel private(@privatization_memref -> %a : memref<10x20xf32>) {
}

OPERATION_NAME = 'acc.private.recipe'

_ODS_REGIONS = (2, True)

sym_name() → _ods_ir

type_() → _ods_ir

initRegion() → _ods_ir

destroyRegion() → _ods_ir

mlir.dialects._acc_ops_gen.private_recipe(sym_name, type_, *, loc=None, ip=None) → PrivateRecipeOp

class mlir.dialects._acc_ops_gen.ReductionOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.reduction'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.reduction(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.ReductionRecipeOp(sym_name, type_, reductionOperator, *, loc=None, ip=None)

Bases: _ods_ir

Declares an OpenACC reduction recipe. The operation requires two mandatory regions and one optional region.

#. The initializer region specifies how to initialize the local reduction value. The region has a first argument that contains the original value that needs to be reduced, followed by bounds arguments (if any) in order from innermost to outermost dimension. It is expected to acc.yield the initialized reduction value. #. The combiner region contains a sequence of operations to combine two values of the reduction type into one. It has the first reduction value, the second reduction value, followed by bounds arguments (if any) in the same order. It is expected to acc.yield the combined value. #. The optional destroy region specifies how to destruct the value when it reaches its end of life. It takes the original value, the reduction value, and bounds arguments (if any) in the same order.

Example:

acc.reduction.recipe @reduction_add_memref : memref<10x20xf32> reduction_operator<add> init {
^bb0(%original: memref<10x20xf32>):
  // init region contains a sequence of operations to initialize the local
  // reduction value as specified in 2.5.15
  %alloca = memref.alloca() : memref<10x20xf32>
  %cst = arith.constant 0.0 : f32
  linalg.fill ins(%cst : f32) outs(%alloca : memref<10x20xf32>)
  acc.yield %alloca : memref<10x20xf32>
} combiner {
^bb0(%lhs: memref<10x20xf32>, %rhs: memref<10x20xf32>):
  // combiner region contains a sequence of operations to combine
  // two values into one.
  linalg.add ins(%lhs, %rhs : memref<10x20xf32>, memref<10x20xf32>)
             outs(%lhs : memref<10x20xf32>)
  acc.yield %lhs : memref<10x20xf32>
} destroy {
^bb0(%original: memref<10x20xf32>, %reduction: memref<10x20xf32>):
  // destroy region is empty since alloca is automatically cleaned up
  acc.terminator
}

// Example with bounds for array slicing:
acc.reduction.recipe @reduction_add_slice : memref<10x20xf32> reduction_operator<add> init {
^bb0(%original: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Extract bounds and create appropriately sized allocation
  %extent_inner = acc.get_extent %bounds_inner : (!acc.data_bounds_ty) -> index
  %extent_outer = acc.get_extent %bounds_outer : (!acc.data_bounds_ty) -> index
  %slice_alloc = memref.alloca(%extent_outer, %extent_inner) : memref<?x?xf32>
  %cst = arith.constant 0.0 : f32
  linalg.fill ins(%cst : f32) outs(%slice_alloc : memref<?x?xf32>)
  // ... base pointer adjustment logic ...
  acc.yield %result : memref<10x20xf32>
} combiner {
^bb0(%lhs: memref<10x20xf32>, %rhs: memref<10x20xf32>, %bounds_inner: !acc.data_bounds_ty, %bounds_outer: !acc.data_bounds_ty):
  // Extract bounds to operate only on the slice portion
  %lb_inner = acc.get_lowerbound %bounds_inner : (!acc.data_bounds_ty) -> index
  %lb_outer = acc.get_lowerbound %bounds_outer : (!acc.data_bounds_ty) -> index
  %extent_inner = acc.get_extent %bounds_inner : (!acc.data_bounds_ty) -> index
  %extent_outer = acc.get_extent %bounds_outer : (!acc.data_bounds_ty) -> index

  // Create subviews to access only the slice portions
  %lhs_slice = memref.subview %lhs[%lb_outer, %lb_inner][%extent_outer, %extent_inner][1, 1]
    : memref<10x20xf32> to memref<?x?xf32, strided<[20, 1], offset: ?>>
  %rhs_slice = memref.subview %rhs[%lb_outer, %lb_inner][%extent_outer, %extent_inner][1, 1]
    : memref<10x20xf32> to memref<?x?xf32, strided<[20, 1], offset: ?>>

  // Combine only the slice portions
  linalg.add ins(%lhs_slice, %rhs_slice : memref<?x?xf32, strided<[20, 1], offset: ?>>, memref<?x?xf32, strided<[20, 1], offset: ?>>)
             outs(%lhs_slice : memref<?x?xf32, strided<[20, 1], offset: ?>>)
  acc.yield %lhs : memref<10x20xf32>
}

// The reduction symbol is then used in the corresponding operation.
acc.parallel reduction(@reduction_add_memref -> %a : memref<10x20xf32>) {
}

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

|------------------------------------------------| | C/C++ | Fortran | |-----------------------|————————| | operator | init value | operator | init value | | + | 0 | + | 0 | | * | 1 | * | 1 | | max | least | max | least | | min | largest | min | largest | | & | ~0 | iand | all bits on | | | | 0 | ior | 0 | | ^ | 0 | ieor | 0 | | && | 1 | .and. | .true. | | || | 0 | .or. | .false. | | | | .eqv. | .true. | | | | .neqv. | .false. | ————————————————-|

OPERATION_NAME = 'acc.reduction.recipe'

_ODS_REGIONS = (3, True)

sym_name() → _ods_ir

type_() → _ods_ir

reductionOperator() → _ods_ir

initRegion() → _ods_ir

combinerRegion() → _ods_ir

destroyRegion() → _ods_ir

mlir.dialects._acc_ops_gen.reduction_recipe(sym_name, type_, reduction_operator, *, loc=None, ip=None) → ReductionRecipeOp

class mlir.dialects._acc_ops_gen.RoutineOp(sym_name, func_name, *, bindIdName=None, bindStrName=None, bindIdNameDeviceType=None, bindStrNameDeviceType=None, worker=None, vector=None, seq=None, nohost=None, implicit=None, gang=None, gangDim=None, gangDimDeviceType=None, loc=None, ip=None)

Bases: _ods_ir

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

bind, gang, worker, vector and seq operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.routine'

_ODS_REGIONS = (0, True)

sym_name() → _ods_ir

func_name() → _ods_ir

bindIdName() → _ods_ir | None

bindStrName() → _ods_ir | None

bindIdNameDeviceType() → _ods_ir | None

bindStrNameDeviceType() → _ods_ir | None

worker() → _ods_ir | None

vector() → _ods_ir | None

seq() → _ods_ir | None

nohost() → bool

implicit() → bool

gang() → _ods_ir | None

gangDim() → _ods_ir | None

gangDimDeviceType() → _ods_ir | None

mlir.dialects._acc_ops_gen.routine(sym_name, func_name, *, bind_id_name=None, bind_str_name=None, bind_id_name_device_type=None, bind_str_name_device_type=None, worker=None, vector=None, seq=None, nohost=None, implicit=None, gang=None, gang_dim=None, gang_dim_device_type=None, loc=None, ip=None) → RoutineOp

class mlir.dialects._acc_ops_gen.SerialOp(asyncOperands, waitOperands, reductionOperands, privateOperands, firstprivateOperands, dataClauseOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, ifCond=None, selfCond=None, selfAttr=None, reductionRecipes=None, privatizationRecipes=None, firstprivatizationRecipes=None, defaultAttr=None, combined=None, loc=None, ip=None)

Bases: _ods_ir

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
}

async and wait operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.serial'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (1, True)

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

ifCond() → _ods_ir | None

selfCond() → _ods_ir | None

reductionOperands() → _ods_ir

privateOperands() → _ods_ir

firstprivateOperands() → _ods_ir

dataClauseOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

selfAttr() → bool

reductionRecipes() → _ods_ir | None

privatizationRecipes() → _ods_ir | None

firstprivatizationRecipes() → _ods_ir | None

defaultAttr() → _ods_ir | None

combined() → bool

region() → _ods_ir

mlir.dialects._acc_ops_gen.serial(async_operands, wait_operands, reduction_operands, private_operands, firstprivate_operands, data_clause_operands, *, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, if_cond=None, self_cond=None, self_attr=None, reduction_recipes=None, privatization_recipes=None, firstprivatization_recipes=None, default_attr=None, combined=None, loc=None, ip=None) → SerialOp

class mlir.dialects._acc_ops_gen.SetOp(*, device_type=None, defaultAsync=None, deviceNum=None, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

acc.set device_num(%dev1 : i32)

OPERATION_NAME = 'acc.set'

_ODS_OPERAND_SEGMENTS = [0, 0, 0]

_ODS_REGIONS = (0, True)

defaultAsync() → _ods_ir | None

deviceNum() → _ods_ir | None

ifCond() → _ods_ir | None

device_type() → _ods_ir | None

mlir.dialects._acc_ops_gen.set(*, device_type=None, default_async=None, device_num=None, if_cond=None, loc=None, ip=None) → SetOp

class mlir.dialects._acc_ops_gen.ShutdownOp(*, device_types=None, deviceNum=None, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

OPERATION_NAME = 'acc.shutdown'

_ODS_OPERAND_SEGMENTS = [0, 0]

_ODS_REGIONS = (0, True)

deviceNum() → _ods_ir | None

ifCond() → _ods_ir | None

device_types() → _ods_ir | None

mlir.dialects._acc_ops_gen.shutdown(*, device_types=None, device_num=None, if_cond=None, loc=None, ip=None) → ShutdownOp

class mlir.dialects._acc_ops_gen.TerminatorOp(*, loc=None, ip=None)

Bases: _ods_ir

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.

OPERATION_NAME = 'acc.terminator'

_ODS_REGIONS = (0, True)

mlir.dialects._acc_ops_gen.terminator(*, loc=None, ip=None) → TerminatorOp

class mlir.dialects._acc_ops_gen.UpdateDeviceOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.update_device'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.update_device(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.UpdateHostOp(accVar, var, varType, bounds, asyncOperands, *, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

• varPtr: The address of variable to copy back to.

• accVar: The acc 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, always, etc) * name: Holds the name of variable as specified in user clause (including bounds). The async values attached to the data exit operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.update_host'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

accVar() → _ods_ir

var() → _ods_ir

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

mlir.dialects._acc_ops_gen.update_host(acc_var, var, var_type, bounds, async_operands, *, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → UpdateHostOp

class mlir.dialects._acc_ops_gen.UpdateOp(asyncOperands, waitOperands, dataClauseOperands, *, ifCond=None, asyncOperandsDeviceType=None, asyncOnly=None, waitOperandsSegments=None, waitOperandsDeviceType=None, hasWaitDevnum=None, waitOnly=None, ifPresent=None, loc=None, ip=None)

Bases: _ods_ir

The acc.update 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}

async and wait operands are supported with device_type information. They should only be accessed by the extra provided getters. If modified, the corresponding device_type attributes must be modified as well.

OPERATION_NAME = 'acc.update'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

ifCond() → _ods_ir | None

asyncOperands() → _ods_ir

waitOperands() → _ods_ir

dataClauseOperands() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

waitOperandsSegments() → _ods_ir | None

waitOperandsDeviceType() → _ods_ir | None

hasWaitDevnum() → _ods_ir | None

waitOnly() → _ods_ir | None

ifPresent() → bool

mlir.dialects._acc_ops_gen.update(async_operands, wait_operands, data_clause_operands, *, if_cond=None, async_operands_device_type=None, async_only=None, wait_operands_segments=None, wait_operands_device_type=None, has_wait_devnum=None, wait_only=None, if_present=None, loc=None, ip=None) → UpdateOp

class mlir.dialects._acc_ops_gen.UseDeviceOp(accVar, var, varType, bounds, asyncOperands, *, varPtrPtr=None, asyncOperandsDeviceType=None, asyncOnly=None, dataClause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None)

Bases: _ods_ir

Description of arguments:

• var: The variable to copy. Must be either MappableType or

PointerLikeType. * varType: The type of the variable that is being copied. When var is a MappableType, this matches the type of var. When var is a PointerLikeType, this type holds information about the target of the pointer. * varPtrPtr: Specifies the address of the address of var - 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. * asyncOperands and asyncOperandsDeviceType: pair-wise lists of the async clause values associated with device_type’s. * asyncOnly: a list of device_type’s for which async clause does not specify a value (default is acc_async_noval - OpenACC 3.3 2.16.1). * 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. * modifiers: Keeps track of the data clause modifiers (eg zero, readonly, etc) * name: Holds the name of variable as specified in user clause (including bounds).

The async values attached to the data entry operation imply that the data action applies to all device types specified by the device_type clauses using the activity queues on these devices as defined by the async values.

OPERATION_NAME = 'acc.use_device'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

var() → _ods_ir

varPtrPtr() → _ods_ir | None

bounds() → _ods_ir

asyncOperands() → _ods_ir

varType() → _ods_ir

asyncOperandsDeviceType() → _ods_ir | None

asyncOnly() → _ods_ir | None

dataClause() → _ods_ir

structured() → _ods_ir

implicit() → _ods_ir

modifiers() → _ods_ir

name() → _ods_ir | None

Returns the fully qualified name of the operation.

accVar() → _ods_ir

mlir.dialects._acc_ops_gen.use_device(acc_var, var, var_type, bounds, async_operands, *, var_ptr_ptr=None, async_operands_device_type=None, async_only=None, data_clause=None, structured=None, implicit=None, modifiers=None, name=None, loc=None, ip=None) → _ods_ir

class mlir.dialects._acc_ops_gen.WaitOp(waitOperands, *, asyncOperand=None, waitDevnum=None, async_=None, ifCond=None, loc=None, ip=None)

Bases: _ods_ir

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

Example:

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

acc.wait does not implement MemoryEffects interface, so it affects all the resources. This is conservatively correct. More precise modelling of the memory effects seems to be impossible without the whole program analysis.

OPERATION_NAME = 'acc.wait'

_ODS_OPERAND_SEGMENTS

_ODS_REGIONS = (0, True)

waitOperands() → _ods_ir

asyncOperand() → _ods_ir | None

waitDevnum() → _ods_ir | None

ifCond() → _ods_ir | None

async_() → bool

mlir.dialects._acc_ops_gen.wait(wait_operands, *, async_operand=None, wait_devnum=None, async_=None, if_cond=None, loc=None, ip=None) → WaitOp

class mlir.dialects._acc_ops_gen.YieldOp(operands_, *, loc=None, ip=None)

Bases: _ods_ir

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.

OPERATION_NAME = 'acc.yield'

_ODS_REGIONS = (0, True)

operands_() → _ods_ir

mlir.dialects._acc_ops_gen.yield_(operands_, *, loc=None, ip=None) → YieldOp