mlir.dialects._acc_ops_gen ========================== .. py:module:: mlir.dialects._acc_ops_gen Attributes ---------- .. autoapisummary:: mlir.dialects._acc_ops_gen._ods_ir mlir.dialects._acc_ops_gen._Buffer Classes ------- .. autoapisummary:: mlir.dialects._acc_ops_gen._Dialect mlir.dialects._acc_ops_gen.AtomicCaptureOp mlir.dialects._acc_ops_gen.AtomicCaptureOpAdaptor mlir.dialects._acc_ops_gen.AtomicReadOp mlir.dialects._acc_ops_gen.AtomicReadOpAdaptor mlir.dialects._acc_ops_gen.AtomicUpdateOp mlir.dialects._acc_ops_gen.AtomicUpdateOpAdaptor mlir.dialects._acc_ops_gen.AtomicWriteOp mlir.dialects._acc_ops_gen.AtomicWriteOpAdaptor mlir.dialects._acc_ops_gen.AttachOp mlir.dialects._acc_ops_gen.AttachOpAdaptor mlir.dialects._acc_ops_gen.CacheOp mlir.dialects._acc_ops_gen.CacheOpAdaptor mlir.dialects._acc_ops_gen.ComputeRegionOp mlir.dialects._acc_ops_gen.ComputeRegionOpAdaptor mlir.dialects._acc_ops_gen.CopyinOp mlir.dialects._acc_ops_gen.CopyinOpAdaptor mlir.dialects._acc_ops_gen.CopyoutOp mlir.dialects._acc_ops_gen.CopyoutOpAdaptor mlir.dialects._acc_ops_gen.CreateOp mlir.dialects._acc_ops_gen.CreateOpAdaptor mlir.dialects._acc_ops_gen.DataBoundsOp mlir.dialects._acc_ops_gen.DataBoundsOpAdaptor mlir.dialects._acc_ops_gen.DataOp mlir.dialects._acc_ops_gen.DataOpAdaptor mlir.dialects._acc_ops_gen.DeclareDeviceResidentOp mlir.dialects._acc_ops_gen.DeclareDeviceResidentOpAdaptor mlir.dialects._acc_ops_gen.DeclareEnterOp mlir.dialects._acc_ops_gen.DeclareEnterOpAdaptor mlir.dialects._acc_ops_gen.DeclareExitOp mlir.dialects._acc_ops_gen.DeclareExitOpAdaptor mlir.dialects._acc_ops_gen.DeclareLinkOp mlir.dialects._acc_ops_gen.DeclareLinkOpAdaptor mlir.dialects._acc_ops_gen.DeclareOp mlir.dialects._acc_ops_gen.DeclareOpAdaptor mlir.dialects._acc_ops_gen.DeleteOp mlir.dialects._acc_ops_gen.DeleteOpAdaptor mlir.dialects._acc_ops_gen.DetachOp mlir.dialects._acc_ops_gen.DetachOpAdaptor mlir.dialects._acc_ops_gen.DevicePtrOp mlir.dialects._acc_ops_gen.DevicePtrOpAdaptor mlir.dialects._acc_ops_gen.EnterDataOp mlir.dialects._acc_ops_gen.EnterDataOpAdaptor mlir.dialects._acc_ops_gen.ExitDataOp mlir.dialects._acc_ops_gen.ExitDataOpAdaptor mlir.dialects._acc_ops_gen.FirstprivateMapInitialOp mlir.dialects._acc_ops_gen.FirstprivateMapInitialOpAdaptor mlir.dialects._acc_ops_gen.FirstprivateOp mlir.dialects._acc_ops_gen.FirstprivateOpAdaptor mlir.dialects._acc_ops_gen.FirstprivateRecipeOp mlir.dialects._acc_ops_gen.FirstprivateRecipeOpAdaptor mlir.dialects._acc_ops_gen.GPUSharedMemoryOp mlir.dialects._acc_ops_gen.GPUSharedMemoryOpAdaptor mlir.dialects._acc_ops_gen.GetDevicePtrOp mlir.dialects._acc_ops_gen.GetDevicePtrOpAdaptor mlir.dialects._acc_ops_gen.GetExtentOp mlir.dialects._acc_ops_gen.GetExtentOpAdaptor mlir.dialects._acc_ops_gen.GetLowerboundOp mlir.dialects._acc_ops_gen.GetLowerboundOpAdaptor mlir.dialects._acc_ops_gen.GetStrideOp mlir.dialects._acc_ops_gen.GetStrideOpAdaptor mlir.dialects._acc_ops_gen.GetUpperboundOp mlir.dialects._acc_ops_gen.GetUpperboundOpAdaptor mlir.dialects._acc_ops_gen.GlobalConstructorOp mlir.dialects._acc_ops_gen.GlobalConstructorOpAdaptor mlir.dialects._acc_ops_gen.GlobalDestructorOp mlir.dialects._acc_ops_gen.GlobalDestructorOpAdaptor mlir.dialects._acc_ops_gen.HostDataOp mlir.dialects._acc_ops_gen.HostDataOpAdaptor mlir.dialects._acc_ops_gen.InitOp mlir.dialects._acc_ops_gen.InitOpAdaptor mlir.dialects._acc_ops_gen.KernelEnvironmentOp mlir.dialects._acc_ops_gen.KernelEnvironmentOpAdaptor mlir.dialects._acc_ops_gen.KernelsOp mlir.dialects._acc_ops_gen.KernelsOpAdaptor mlir.dialects._acc_ops_gen.LoopOp mlir.dialects._acc_ops_gen.LoopOpAdaptor mlir.dialects._acc_ops_gen.MapInfoOp mlir.dialects._acc_ops_gen.MapInfoOpAdaptor mlir.dialects._acc_ops_gen.NoCreateOp mlir.dialects._acc_ops_gen.NoCreateOpAdaptor mlir.dialects._acc_ops_gen.OnDeviceOp mlir.dialects._acc_ops_gen.OnDeviceOpAdaptor mlir.dialects._acc_ops_gen.ParWidthOp mlir.dialects._acc_ops_gen.ParWidthOpAdaptor mlir.dialects._acc_ops_gen.ParallelOp mlir.dialects._acc_ops_gen.ParallelOpAdaptor mlir.dialects._acc_ops_gen.PredicateRegionOp mlir.dialects._acc_ops_gen.PredicateRegionOpAdaptor mlir.dialects._acc_ops_gen.PresentOp mlir.dialects._acc_ops_gen.PresentOpAdaptor mlir.dialects._acc_ops_gen.PrivateLocalOp mlir.dialects._acc_ops_gen.PrivateLocalOpAdaptor mlir.dialects._acc_ops_gen.PrivateOp mlir.dialects._acc_ops_gen.PrivateOpAdaptor mlir.dialects._acc_ops_gen.PrivateRecipeOp mlir.dialects._acc_ops_gen.PrivateRecipeOpAdaptor mlir.dialects._acc_ops_gen.PrivatizeOp mlir.dialects._acc_ops_gen.PrivatizeOpAdaptor mlir.dialects._acc_ops_gen.ReductionAccumulateArrayOp mlir.dialects._acc_ops_gen.ReductionAccumulateArrayOpAdaptor mlir.dialects._acc_ops_gen.ReductionAccumulateOp mlir.dialects._acc_ops_gen.ReductionAccumulateOpAdaptor mlir.dialects._acc_ops_gen.ReductionCombineOp mlir.dialects._acc_ops_gen.ReductionCombineOpAdaptor mlir.dialects._acc_ops_gen.ReductionCombineRegionOp mlir.dialects._acc_ops_gen.ReductionCombineRegionOpAdaptor mlir.dialects._acc_ops_gen.ReductionInitOp mlir.dialects._acc_ops_gen.ReductionInitOpAdaptor mlir.dialects._acc_ops_gen.ReductionOp mlir.dialects._acc_ops_gen.ReductionOpAdaptor mlir.dialects._acc_ops_gen.ReductionRecipeOp mlir.dialects._acc_ops_gen.ReductionRecipeOpAdaptor mlir.dialects._acc_ops_gen.RoutineOp mlir.dialects._acc_ops_gen.RoutineOpAdaptor mlir.dialects._acc_ops_gen.SerialOp mlir.dialects._acc_ops_gen.SerialOpAdaptor mlir.dialects._acc_ops_gen.SetOp mlir.dialects._acc_ops_gen.SetOpAdaptor mlir.dialects._acc_ops_gen.ShutdownOp mlir.dialects._acc_ops_gen.ShutdownOpAdaptor mlir.dialects._acc_ops_gen.TerminatorOp mlir.dialects._acc_ops_gen.TerminatorOpAdaptor mlir.dialects._acc_ops_gen.UnwrapPrivateOp mlir.dialects._acc_ops_gen.UnwrapPrivateOpAdaptor mlir.dialects._acc_ops_gen.UpdateDeviceOp mlir.dialects._acc_ops_gen.UpdateDeviceOpAdaptor mlir.dialects._acc_ops_gen.UpdateHostOp mlir.dialects._acc_ops_gen.UpdateHostOpAdaptor mlir.dialects._acc_ops_gen.UpdateOp mlir.dialects._acc_ops_gen.UpdateOpAdaptor mlir.dialects._acc_ops_gen.UseDeviceOp mlir.dialects._acc_ops_gen.UseDeviceOpAdaptor mlir.dialects._acc_ops_gen.WaitOp mlir.dialects._acc_ops_gen.WaitOpAdaptor mlir.dialects._acc_ops_gen.YieldOp mlir.dialects._acc_ops_gen.YieldOpAdaptor Functions --------- .. autoapisummary:: mlir.dialects._acc_ops_gen.atomic_capture mlir.dialects._acc_ops_gen.atomic_read mlir.dialects._acc_ops_gen.atomic_update mlir.dialects._acc_ops_gen.atomic_write mlir.dialects._acc_ops_gen.attach mlir.dialects._acc_ops_gen.cache mlir.dialects._acc_ops_gen.compute_region mlir.dialects._acc_ops_gen.copyin mlir.dialects._acc_ops_gen.copyout mlir.dialects._acc_ops_gen.create_ mlir.dialects._acc_ops_gen.bounds mlir.dialects._acc_ops_gen.data mlir.dialects._acc_ops_gen.declare_device_resident mlir.dialects._acc_ops_gen.declare_enter mlir.dialects._acc_ops_gen.declare_exit mlir.dialects._acc_ops_gen.declare_link mlir.dialects._acc_ops_gen.declare mlir.dialects._acc_ops_gen.delete mlir.dialects._acc_ops_gen.detach mlir.dialects._acc_ops_gen.deviceptr mlir.dialects._acc_ops_gen.enter_data mlir.dialects._acc_ops_gen.exit_data mlir.dialects._acc_ops_gen.firstprivate_map mlir.dialects._acc_ops_gen.firstprivate mlir.dialects._acc_ops_gen.firstprivate_recipe mlir.dialects._acc_ops_gen.gpu_shared_memory mlir.dialects._acc_ops_gen.getdeviceptr mlir.dialects._acc_ops_gen.get_extent mlir.dialects._acc_ops_gen.get_lowerbound mlir.dialects._acc_ops_gen.get_stride mlir.dialects._acc_ops_gen.get_upperbound mlir.dialects._acc_ops_gen.global_ctor mlir.dialects._acc_ops_gen.global_dtor mlir.dialects._acc_ops_gen.host_data mlir.dialects._acc_ops_gen.init mlir.dialects._acc_ops_gen.kernel_environment mlir.dialects._acc_ops_gen.kernels mlir.dialects._acc_ops_gen.loop mlir.dialects._acc_ops_gen.map_info mlir.dialects._acc_ops_gen.nocreate mlir.dialects._acc_ops_gen.on_device mlir.dialects._acc_ops_gen.par_width mlir.dialects._acc_ops_gen.parallel mlir.dialects._acc_ops_gen.predicate_region mlir.dialects._acc_ops_gen.present mlir.dialects._acc_ops_gen.private_local mlir.dialects._acc_ops_gen.private mlir.dialects._acc_ops_gen.private_recipe mlir.dialects._acc_ops_gen.privatize mlir.dialects._acc_ops_gen.reduction_accumulate_array mlir.dialects._acc_ops_gen.reduction_accumulate mlir.dialects._acc_ops_gen.reduction_combine mlir.dialects._acc_ops_gen.reduction_combine_region mlir.dialects._acc_ops_gen.reduction_init mlir.dialects._acc_ops_gen.reduction mlir.dialects._acc_ops_gen.reduction_recipe mlir.dialects._acc_ops_gen.routine mlir.dialects._acc_ops_gen.serial mlir.dialects._acc_ops_gen.set mlir.dialects._acc_ops_gen.shutdown mlir.dialects._acc_ops_gen.terminator mlir.dialects._acc_ops_gen.unwrap_private mlir.dialects._acc_ops_gen.update_device mlir.dialects._acc_ops_gen.update_host mlir.dialects._acc_ops_gen.update mlir.dialects._acc_ops_gen.use_device mlir.dialects._acc_ops_gen.wait mlir.dialects._acc_ops_gen.yield_ Module Contents --------------- .. py:data:: _ods_ir .. py:data:: _Buffer .. py:class:: _Dialect(descriptor: object) Bases: :py:obj:`_ods_ir` .. py:attribute:: DIALECT_NAMESPACE :value: 'acc' .. py:class:: AtomicCaptureOp(*, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` This operation performs an atomic capture. The region has the following allowed forms: .. code:: 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 } .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.capture' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: region() -> _ods_ir .. py:class:: AtomicCaptureOpAdaptor(operands: list[Value], attributes: OpAttributeMap) AtomicCaptureOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.capture' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:function:: atomic_capture(*, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> AtomicCaptureOp .. py:class:: AtomicReadOp(x: _ods_ir, v: _ods_ir, element_type: Union[_ods_ir, _ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.read' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: x() -> _ods_ir .. py:method:: v() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: element_type() -> _ods_ir .. py:class:: AtomicReadOpAdaptor(operands: list[Value], attributes: OpAttributeMap) AtomicReadOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.read' .. py:method:: x() -> _ods_ir .. py:method:: v() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: element_type() -> _ods_ir .. py:function:: atomic_read(x: _ods_ir, v: _ods_ir, element_type: Union[_ods_ir, _ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> AtomicReadOp .. py:class:: AtomicUpdateOp(x: _ods_ir, *, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.update' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: x() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: region() -> _ods_ir .. py:class:: AtomicUpdateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) AtomicUpdateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.update' .. py:method:: x() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:function:: atomic_update(x: _ods_ir, *, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> AtomicUpdateOp .. py:class:: AtomicWriteOp(x: _ods_ir, expr: _ods_ir, *, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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). .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.write' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: x() -> _ods_ir .. py:method:: expr() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:class:: AtomicWriteOpAdaptor(operands: list[Value], attributes: OpAttributeMap) AtomicWriteOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.atomic.write' .. py:method:: x() -> _ods_ir .. py:method:: expr() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:function:: atomic_write(x: _ods_ir, expr: _ods_ir, *, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> AtomicWriteOp .. py:class:: AttachOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.attach' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: AttachOpAdaptor(operands: list[Value], attributes: OpAttributeMap) AttachOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.attach' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: attach(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: CacheOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.cache' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: CacheOpAdaptor(operands: list[Value], attributes: OpAttributeMap) CacheOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.cache' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: cache(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: ComputeRegionOp(results_: Sequence[_ods_ir], launchArgs: Sequence[_ods_ir[_ods_ir]], inputArgs: Sequence[_ods_ir], origin: Union[str, _ods_ir], *, stream: Optional[_ods_ir] = None, kernel_func_name: Optional[Union[str, _ods_ir]] = None, kernel_module_name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The ``acc.compute_region`` operation wraps a region of code that will be compiled and executed on a GPU. It is typically produced by lowering OpenACC compute constructs (``acc.parallel``, ``acc.kernels``, ``acc.serial``) but can also be targeted directly by other frontends or lowered from other constructs that benefit from the automatic parallelization and data mapping facilities that the ``acc`` dialect provides. It serves as the bridge between the high-level representation and the ``gpu.launch`` operation. The operation is ``IsolatedFromAbove``: all values used inside the region must be explicitly captured. Values are captured in two ways: * Launch arguments (``launch``): Results of ``acc.par_width`` operations that define the parallel launch configuration. These become ``index``-typed block arguments representing the parallel width for each dimension. * Input arguments (``ins``): Arbitrary values captured from outside the region (data pointers, scalars, etc.). These become block arguments with their original types. The ``origin`` attribute records which construct produced this compute region (e.g., ``"acc.parallel"``, ``"acc.kernels"``). This is intended to be solely informational. Canonicalization may simplify ``ins`` captures: duplicate ``ins`` operands (same SSA value threaded more than once) are merged by reusing the first block argument, and unused ``ins`` operands (block arguments with no uses) are removed. ``launch`` operands are never merged or dropped. Example: .. code:: mlir %w0 = acc.par_width %c128 par_dim(#acc.par_dim) %w1 = acc.par_width %c8 par_dim(#acc.par_dim) acc.compute_region launch(%arg0 = %w0, %arg1 = %w1) ins(%arg2 = %data) : (memref<1024xf32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c1024 = arith.constant 1024 : index scf.parallel (%iv) = (%c0) to (%c1024) step (%c1) { %v = memref.load %arg2[%iv] : memref<1024xf32> scf.reduce } {acc.par_dims = #acc} acc.yield } <{origin = "acc.parallel"}> .. py:attribute:: OPERATION_NAME :value: 'acc.compute_region' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: launchArgs() -> _ods_ir[_ods_ir] .. py:method:: inputArgs() -> _ods_ir .. py:method:: stream() -> Optional[_ods_ir] .. py:method:: origin() -> _ods_ir .. py:method:: kernel_func_name() -> Optional[_ods_ir] .. py:method:: kernel_module_name() -> Optional[_ods_ir] .. py:method:: results_() -> _ods_ir .. py:method:: region() -> _ods_ir .. py:class:: ComputeRegionOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ComputeRegionOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.compute_region' .. py:method:: launchArgs() -> _ods_ir[_ods_ir] .. py:method:: inputArgs() -> _ods_ir .. py:method:: stream() -> Optional[_ods_ir] .. py:method:: origin() -> _ods_ir .. py:method:: kernel_func_name() -> Optional[_ods_ir] .. py:method:: kernel_module_name() -> Optional[_ods_ir] .. py:function:: compute_region(results_: Sequence[_ods_ir], launch_args: Sequence[_ods_ir[_ods_ir]], input_args: Sequence[_ods_ir], origin: Union[str, _ods_ir], *, stream: Optional[_ods_ir] = None, kernel_func_name: Optional[Union[str, _ods_ir]] = None, kernel_module_name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> Union[_ods_ir, _ods_ir, ComputeRegionOp] .. py:class:: CopyinOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.copyin' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: CopyinOpAdaptor(operands: list[Value], attributes: OpAttributeMap) CopyinOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.copyin' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: copyin(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: CopyoutOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.copyout' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: accVar() -> _ods_ir .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:class:: CopyoutOpAdaptor(operands: list[Value], attributes: OpAttributeMap) CopyoutOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.copyout' .. py:method:: accVar() -> _ods_ir .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:function:: copyout(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> CopyoutOp .. py:class:: CreateOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.create' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: CreateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) CreateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.create' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: create_(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: DataBoundsOp(*, lowerbound: Optional[_ods_ir] = None, upperbound: Optional[_ods_ir] = None, extent: Optional[_ods_ir] = None, sourceExtent: Optional[_ods_ir] = None, stride: Optional[_ods_ir] = None, strideInBytes: Optional[Union[bool, _ods_ir]] = None, startIdx: Optional[_ods_ir] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. ``sourceExtent`` is the extent of the corresponding dimension in the original array. It may differ from ``extent``, which describes the selected section. When absent, the source extent is the same as ``extent``. 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++: .. code:: int array[10]; #pragma acc copy(array[1:9]) => .. code:: mlir acc.bounds lb(1) ub(9) extent(9) startIdx(0) stride(1) Fortran: .. code:: integer :: array(1:10) !$acc copy(array(2:10)) => .. code:: mlir acc.bounds lb(1) ub(9) extent(9) startIdx(1) stride(1) .. py:attribute:: OPERATION_NAME :value: 'acc.bounds' .. py:attribute:: _ODS_OPERAND_SEGMENTS :value: [0, 0, 0, 0, 0, 0] .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: lowerbound() -> Optional[_ods_ir] .. py:method:: upperbound() -> Optional[_ods_ir] .. py:method:: extent() -> Optional[_ods_ir] .. py:method:: sourceExtent() -> Optional[_ods_ir] .. py:method:: stride() -> Optional[_ods_ir] .. py:method:: startIdx() -> Optional[_ods_ir] .. py:method:: strideInBytes() -> _ods_ir .. py:method:: result() -> _ods_ir Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: DataBoundsOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DataBoundsOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.bounds' .. py:method:: lowerbound() -> Optional[_ods_ir] .. py:method:: upperbound() -> Optional[_ods_ir] .. py:method:: extent() -> Optional[_ods_ir] .. py:method:: sourceExtent() -> Optional[_ods_ir] .. py:method:: stride() -> Optional[_ods_ir] .. py:method:: startIdx() -> Optional[_ods_ir] .. py:method:: strideInBytes() -> _ods_ir .. py:function:: bounds(*, lowerbound: Optional[_ods_ir] = None, upperbound: Optional[_ods_ir] = None, extent: Optional[_ods_ir] = None, source_extent: Optional[_ods_ir] = None, stride: Optional[_ods_ir] = None, stride_in_bytes: Optional[Union[bool, _ods_ir]] = None, start_idx: Optional[_ods_ir] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: DataOp(asyncOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, waitOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, waitOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, hasWaitDevnum: Optional[Union[Sequence[bool], _ods_ir]] = None, waitOnly: Optional[Union[Any, _ods_ir]] = None, defaultAttr: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.data' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: region() -> _ods_ir .. py:class:: DataOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DataOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.data' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:function:: data(async_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, wait_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, wait_operands_device_type: Optional[Union[Any, _ods_ir]] = None, has_wait_devnum: Optional[Union[Sequence[bool], _ods_ir]] = None, wait_only: Optional[Union[Any, _ods_ir]] = None, default_attr: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> DataOp .. py:class:: DeclareDeviceResidentOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.declare_device_resident' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: DeclareDeviceResidentOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeclareDeviceResidentOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.declare_device_resident' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: declare_device_resident(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: DeclareEnterOp(dataClauseOperands: Sequence[_ods_ir], *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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)``: .. code:: mlir %0 = acc.create varPtr(%a : !llvm.ptr) -> !llvm.ptr acc.declare_enter dataOperands(%0 : !llvm.ptr) .. py:attribute:: OPERATION_NAME :value: 'acc.declare_enter' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: token() -> _ods_ir .. py:class:: DeclareEnterOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeclareEnterOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.declare_enter' .. py:method:: dataClauseOperands() -> _ods_ir .. py:function:: declare_enter(data_clause_operands: Sequence[_ods_ir], *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: DeclareExitOp(dataClauseOperands: Sequence[_ods_ir], *, token: Optional[_ods_ir] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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)``: .. code:: mlir %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr {dataClause = #acc.data_clause} acc.declare_exit dataOperands(%0 : !llvm.ptr) acc.delete accPtr(%0 : !llvm.ptr) {dataClause = #acc.data_clause} .. py:attribute:: OPERATION_NAME :value: 'acc.declare_exit' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: token() -> Optional[_ods_ir] .. py:method:: dataClauseOperands() -> _ods_ir .. py:class:: DeclareExitOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeclareExitOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.declare_exit' .. py:method:: token() -> Optional[_ods_ir] .. py:method:: dataClauseOperands() -> _ods_ir .. py:function:: declare_exit(data_clause_operands: Sequence[_ods_ir], *, token: Optional[_ods_ir] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> DeclareExitOp .. py:class:: DeclareLinkOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.declare_link' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: DeclareLinkOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeclareLinkOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.declare_link' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: declare_link(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: DeclareOp(dataClauseOperands: Sequence[_ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.declare" operation represents an implicit declare region in function (and subroutine in Fortran). Example: .. code:: mlir %pa = acc.present varPtr(%a : memref<10x10xf32>) -> memref<10x10xf32> acc.declare dataOperands(%pa: memref<10x10xf32>) { // implicit region } .. py:attribute:: OPERATION_NAME :value: 'acc.declare' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: region() -> _ods_ir .. py:class:: DeclareOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeclareOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.declare' .. py:method:: dataClauseOperands() -> _ods_ir .. py:function:: declare(data_clause_operands: Sequence[_ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> DeclareOp .. py:class:: DeleteOp(accVar: _ods_ir, bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.delete' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: accVar() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:class:: DeleteOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DeleteOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.delete' .. py:method:: accVar() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:function:: delete(acc_var: _ods_ir, bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> DeleteOp .. py:class:: DetachOp(accVar: _ods_ir, bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.detach' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: accVar() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:class:: DetachOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DetachOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.detach' .. py:method:: accVar() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:function:: detach(acc_var: _ods_ir, bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> DetachOp .. py:class:: DevicePtrOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.deviceptr' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: DevicePtrOpAdaptor(operands: list[Value], attributes: OpAttributeMap) DevicePtrOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.deviceptr' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: deviceptr(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: EnterDataOp(waitOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, asyncOperand: Optional[_ods_ir] = None, async_: Optional[bool] = None, waitDevnum: Optional[_ods_ir] = None, wait: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.enter_data" operation represents the OpenACC enter data directive. Example: .. code:: mlir acc.enter_data create(%d1 : memref<10xf32>) attributes {async} .. py:attribute:: OPERATION_NAME :value: 'acc.enter_data' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: async_() -> bool .. py:method:: wait() -> bool .. py:class:: EnterDataOpAdaptor(operands: list[Value], attributes: OpAttributeMap) EnterDataOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.enter_data' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: async_() -> bool .. py:method:: wait() -> bool .. py:function:: enter_data(wait_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, async_operand: Optional[_ods_ir] = None, async_: Optional[bool] = None, wait_devnum: Optional[_ods_ir] = None, wait: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> EnterDataOp .. py:class:: ExitDataOp(waitOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, asyncOperand: Optional[_ods_ir] = None, async_: Optional[bool] = None, waitDevnum: Optional[_ods_ir] = None, wait: Optional[bool] = None, finalize: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.exit_data" operation represents the OpenACC exit data directive. Example: .. code:: mlir acc.exit_data delete(%d1 : memref<10xf32>) attributes {async} .. py:attribute:: OPERATION_NAME :value: 'acc.exit_data' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: async_() -> bool .. py:method:: wait() -> bool .. py:method:: finalize() -> bool .. py:class:: ExitDataOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ExitDataOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.exit_data' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: async_() -> bool .. py:method:: wait() -> bool .. py:method:: finalize() -> bool .. py:function:: exit_data(wait_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, async_operand: Optional[_ods_ir] = None, async_: Optional[bool] = None, wait_devnum: Optional[_ods_ir] = None, wait: Optional[bool] = None, finalize: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ExitDataOp .. py:class:: FirstprivateMapInitialOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The ``acc.firstprivate_map`` operation is an intermediate representation used during the decomposition of ``acc.firstprivate`` operations. It represents the mapping of the initial value from the host to the device, which is then used to initialize per-thread private copies. This operation is distinct from ``acc.copyin`` because: * ``acc.copyin`` includes present counter updates, but private variables do not impact reference counters * The mapped value is used to initialize private copies rather than being accessed directly .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate_map' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: FirstprivateMapInitialOpAdaptor(operands: list[Value], attributes: OpAttributeMap) FirstprivateMapInitialOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate_map' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: firstprivate_map(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: FirstprivateOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: FirstprivateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) FirstprivateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: firstprivate(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: FirstprivateRecipeOp(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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 first and may yield additional values that are used only for destruction. #. 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, any additional destruction values yielded by the init region, 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: .. code:: mlir 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 // ... 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> // Copy subview to privatized... acc.terminator } // The privatization symbol is then used in the corresponding operation. acc.parallel firstprivate(@firstprivate_memref -> %a : memref<10x20xf32>) { } .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate.recipe' .. py:attribute:: _ODS_REGIONS :value: (3, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:method:: initRegion() -> _ods_ir .. py:method:: copyRegion() -> _ods_ir .. py:method:: destroyRegion() -> _ods_ir .. py:class:: FirstprivateRecipeOpAdaptor(operands: list[Value], attributes: OpAttributeMap) FirstprivateRecipeOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.firstprivate.recipe' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:function:: firstprivate_recipe(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> FirstprivateRecipeOp .. py:class:: GPUSharedMemoryOp(result: _ods_ir, num_copies: Union[int, _ods_ir], static_upper_bound_bytes: Union[int, _ods_ir], dynamic_sizes: Sequence[_ods_ir[_ods_ir]], *, dynamic_shared_memory_scaling_bytes: Optional[Union[int, _ods_ir]] = None, dynamic_shared_memory_fixed_bytes: Optional[Union[int, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Represents a GPU workgroup-memory allocation in a compute region. The result is a typed ``memref`` view into a byte slab which is later replaced by ``memref.view`` into a dynamic shared-memory blob at the byte offset within the workgroup allocation. Each operation occupies a distinct slot in that collective allocation. ``static_upper_bound_bytes`` is the conservative byte-size upper bound for the slot. ``dynamic_sizes`` supply values for dynamic memref result dimensions. The in-kernel layout of the slot is given by ``static_upper_bound_bytes`` and ``dynamic_sizes``. Optional ``dynamic_shared_memory_scaling_bytes`` and ``dynamic_shared_memory_fixed_bytes`` parameterize ``dynamic_shared_memory_size`` when the slot footprint depends on launch geometry. They must be specified together. When present: dynamic_shared_memory_size = dynamic_shared_memory_scaling_bytes * W + dynamic_shared_memory_fixed_bytes where ``W`` is the launch width that scales the allocation. This linear model arises for ``acc.cache`` regions with dynamic bounds: one cache dimension grows with the thread-parallel launch width, while ``dynamic_shared_memory_fixed_bytes`` covers bytes that do not scale (for example overlap cells at cache tile boundaries so threads can read neighboring source elements without extra global memory traffic). For a 1D dynamic cache with stencil extent ``E``: dynamic_shared_memory_fixed_bytes = (E - 1) * dynamic_shared_memory_scaling_bytes .. code:: Global array: ... | a | b | c | d | e | f | ... [---- cached tile ----] Thread 0 primary: a (reads neighbor b) Thread 1 primary: b (reads neighbors a, c) ... Scaling portion: dynamic_shared_memory_scaling_bytes * W Fixed portion: dynamic_shared_memory_fixed_bytes ((E - 1) cells) The scaling attributes affect only ``dynamic_shared_memory_size``, not the slot layout. For purely static allocations both are omitted and ``dynamic_shared_memory_size`` is the sum of aligned ``static_upper_bound_bytes`` across all slots. Example: .. code:: mlir %sz = arith.constant 128 : index %cache = acc.gpu_shared_memory(%sz) <{num_copies = 1 : i64, static_upper_bound_bytes = 1560 : i64, dynamic_shared_memory_scaling_bytes = 12 : i64, dynamic_shared_memory_fixed_bytes = 24 : i64}> : (index) -> memref> .. py:attribute:: OPERATION_NAME :value: 'acc.gpu_shared_memory' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: dynamic_sizes() -> _ods_ir[_ods_ir] .. py:method:: num_copies() -> _ods_ir .. py:method:: static_upper_bound_bytes() -> _ods_ir .. py:method:: dynamic_shared_memory_scaling_bytes() -> Optional[_ods_ir] .. py:method:: dynamic_shared_memory_fixed_bytes() -> Optional[_ods_ir] .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: GPUSharedMemoryOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GPUSharedMemoryOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.gpu_shared_memory' .. py:method:: dynamic_sizes() -> _ods_ir[_ods_ir] .. py:method:: num_copies() -> _ods_ir .. py:method:: static_upper_bound_bytes() -> _ods_ir .. py:method:: dynamic_shared_memory_scaling_bytes() -> Optional[_ods_ir] .. py:method:: dynamic_shared_memory_fixed_bytes() -> Optional[_ods_ir] .. py:function:: gpu_shared_memory(result: _ods_ir, num_copies: Union[int, _ods_ir], static_upper_bound_bytes: Union[int, _ods_ir], dynamic_sizes: Sequence[_ods_ir[_ods_ir]], *, dynamic_shared_memory_scaling_bytes: Optional[Union[int, _ods_ir]] = None, dynamic_shared_memory_fixed_bytes: Optional[Union[int, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: GetDevicePtrOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.getdeviceptr' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: GetDevicePtrOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GetDevicePtrOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.getdeviceptr' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: getdeviceptr(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: GetExtentOp(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir %extent = acc.get_extent %bounds : (!acc.data_bounds_ty) -> index .. py:attribute:: OPERATION_NAME :value: 'acc.get_extent' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: bounds() -> _ods_ir .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: GetExtentOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GetExtentOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.get_extent' .. py:method:: bounds() -> _ods_ir .. py:function:: get_extent(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: GetLowerboundOp(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir %lb = acc.get_lowerbound %bounds : (!acc.data_bounds_ty) -> index .. py:attribute:: OPERATION_NAME :value: 'acc.get_lowerbound' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: bounds() -> _ods_ir .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: GetLowerboundOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GetLowerboundOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.get_lowerbound' .. py:method:: bounds() -> _ods_ir .. py:function:: get_lowerbound(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: GetStrideOp(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir %stride = acc.get_stride %bounds : (!acc.data_bounds_ty) -> index .. py:attribute:: OPERATION_NAME :value: 'acc.get_stride' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: bounds() -> _ods_ir .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: GetStrideOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GetStrideOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.get_stride' .. py:method:: bounds() -> _ods_ir .. py:function:: get_stride(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: GetUpperboundOp(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir %ub = acc.get_upperbound %bounds : (!acc.data_bounds_ty) -> index .. py:attribute:: OPERATION_NAME :value: 'acc.get_upperbound' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: bounds() -> _ods_ir .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: GetUpperboundOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GetUpperboundOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.get_upperbound' .. py:method:: bounds() -> _ods_ir .. py:function:: get_upperbound(bounds: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: GlobalConstructorOp(sym_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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) } .. py:attribute:: OPERATION_NAME :value: 'acc.global_ctor' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: region() -> _ods_ir .. py:class:: GlobalConstructorOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GlobalConstructorOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.global_ctor' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:function:: global_ctor(sym_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> GlobalConstructorOp .. py:class:: GlobalDestructorOp(sym_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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} acc.declare_exit dataOperands(%1 : !llvm.ptr) acc.delete accPtr(%1 : !llvm.ptr) {dataClause = #acc.data_clause} } .. py:attribute:: OPERATION_NAME :value: 'acc.global_dtor' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: region() -> _ods_ir .. py:class:: GlobalDestructorOpAdaptor(operands: list[Value], attributes: OpAttributeMap) GlobalDestructorOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.global_dtor' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:function:: global_dtor(sym_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> GlobalDestructorOp .. py:class:: HostDataOp(dataClauseOperands: Sequence[_ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, ifPresent: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.host_data" operation represents the OpenACC host_data construct. Example: .. code:: mlir %0 = acc.use_device varPtr(%a : !llvm.ptr) -> !llvm.ptr acc.host_data dataOperands(%0 : !llvm.ptr) { } .. py:attribute:: OPERATION_NAME :value: 'acc.host_data' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: ifPresent() -> bool .. py:method:: region() -> _ods_ir .. py:class:: HostDataOpAdaptor(operands: list[Value], attributes: OpAttributeMap) HostDataOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.host_data' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: ifPresent() -> bool .. py:function:: host_data(data_clause_operands: Sequence[_ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, if_present: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> HostDataOp .. py:class:: InitOp(*, device_types: Optional[Union[Any, _ods_ir]] = None, deviceNum: Optional[_ods_ir] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.init" operation represents the OpenACC init executable directive. Example: .. code:: mlir acc.init acc.init device_num(%dev1 : i32) .. py:attribute:: OPERATION_NAME :value: 'acc.init' .. py:attribute:: _ODS_OPERAND_SEGMENTS :value: [0, 0] .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_types() -> Optional[_ods_ir] .. py:class:: InitOpAdaptor(operands: list[Value], attributes: OpAttributeMap) InitOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.init' .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_types() -> Optional[_ods_ir] .. py:function:: init(*, device_types: Optional[Union[Any, _ods_ir]] = None, device_num: Optional[_ods_ir] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> InitOp .. py:class:: KernelEnvironmentOp(dataClauseOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], *, asyncOperand: Optional[_ods_ir] = None, asyncOnly: Optional[bool] = None, waitDevnum: Optional[_ods_ir] = None, waitOnly: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.kernel_environment' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: asyncOnly() -> bool .. py:method:: waitOnly() -> bool .. py:method:: region() -> _ods_ir .. py:class:: KernelEnvironmentOpAdaptor(operands: list[Value], attributes: OpAttributeMap) KernelEnvironmentOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.kernel_environment' .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: waitOperands() -> _ods_ir .. py:method:: asyncOnly() -> bool .. py:method:: waitOnly() -> bool .. py:function:: kernel_environment(data_clause_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], *, async_operand: Optional[_ods_ir] = None, async_only: Optional[bool] = None, wait_devnum: Optional[_ods_ir] = None, wait_only: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> KernelEnvironmentOp .. py:class:: KernelsOp(asyncOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], numGangs: Sequence[_ods_ir], numWorkers: Sequence[_ods_ir], vectorLength: Sequence[_ods_ir], reductionOperands: Sequence[_ods_ir], privateOperands: Sequence[_ods_ir], firstprivateOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, waitOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, waitOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, hasWaitDevnum: Optional[Union[Sequence[bool], _ods_ir]] = None, waitOnly: Optional[Union[Any, _ods_ir]] = None, numGangsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, numGangsDeviceType: Optional[Union[Any, _ods_ir]] = None, numWorkersDeviceType: Optional[Union[Any, _ods_ir]] = None, vectorLengthDeviceType: Optional[Union[Any, _ods_ir]] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, selfCond: Optional[_ods_ir[_ods_ir]] = None, selfAttr: Optional[bool] = None, defaultAttr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.kernels' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: numGangs() -> _ods_ir .. py:method:: numWorkers() -> _ods_ir .. py:method:: vectorLength() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: numGangsSegments() -> Optional[_ods_ir] .. py:method:: numGangsDeviceType() -> Optional[_ods_ir] .. py:method:: numWorkersDeviceType() -> Optional[_ods_ir] .. py:method:: vectorLengthDeviceType() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:method:: region() -> _ods_ir .. py:class:: KernelsOpAdaptor(operands: list[Value], attributes: OpAttributeMap) KernelsOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.kernels' .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: numGangs() -> _ods_ir .. py:method:: numWorkers() -> _ods_ir .. py:method:: vectorLength() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: numGangsSegments() -> Optional[_ods_ir] .. py:method:: numGangsDeviceType() -> Optional[_ods_ir] .. py:method:: numWorkersDeviceType() -> Optional[_ods_ir] .. py:method:: vectorLengthDeviceType() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:function:: kernels(async_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], num_gangs: Sequence[_ods_ir], num_workers: Sequence[_ods_ir], vector_length: Sequence[_ods_ir], reduction_operands: Sequence[_ods_ir], private_operands: Sequence[_ods_ir], firstprivate_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, wait_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, wait_operands_device_type: Optional[Union[Any, _ods_ir]] = None, has_wait_devnum: Optional[Union[Sequence[bool], _ods_ir]] = None, wait_only: Optional[Union[Any, _ods_ir]] = None, num_gangs_segments: Optional[Union[Sequence[int], _ods_ir]] = None, num_gangs_device_type: Optional[Union[Any, _ods_ir]] = None, num_workers_device_type: Optional[Union[Any, _ods_ir]] = None, vector_length_device_type: Optional[Union[Any, _ods_ir]] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, self_cond: Optional[_ods_ir[_ods_ir]] = None, self_attr: Optional[bool] = None, default_attr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> KernelsOp .. py:class:: LoopOp(results_: Sequence[_ods_ir], lowerbound: Sequence[_ods_ir], upperbound: Sequence[_ods_ir], step: Sequence[_ods_ir], gangOperands: Sequence[_ods_ir], workerNumOperands: Sequence[_ods_ir], vectorOperands: Sequence[_ods_ir], tileOperands: Sequence[_ods_ir], cacheOperands: Sequence[_ods_ir], privateOperands: Sequence[_ods_ir], firstprivateOperands: Sequence[_ods_ir], reductionOperands: Sequence[_ods_ir], *, inclusiveUpperbound: Optional[Union[Sequence[bool], _ods_ir]] = None, collapse: Optional[Union[Sequence[int], _ods_ir]] = None, collapseDeviceType: Optional[Union[Any, _ods_ir]] = None, gangOperandsArgType: Optional[Union[Any, _ods_ir]] = None, gangOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, gangOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, workerNumOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, vectorOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, seq: Optional[Union[Any, _ods_ir]] = None, independent: Optional[Union[Any, _ods_ir]] = None, auto_: Optional[Union[Any, _ods_ir]] = None, gang: Optional[Union[Any, _ods_ir]] = None, worker: Optional[Union[Any, _ods_ir]] = None, vector: Optional[Union[Any, _ods_ir]] = None, tileOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, tileOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, combined: Optional[Union[Any, _ods_ir]] = None, unstructured: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.loop' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: lowerbound() -> _ods_ir .. py:method:: upperbound() -> _ods_ir .. py:method:: step() -> _ods_ir .. py:method:: gangOperands() -> _ods_ir .. py:method:: workerNumOperands() -> _ods_ir .. py:method:: vectorOperands() -> _ods_ir .. py:method:: tileOperands() -> _ods_ir .. py:method:: cacheOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: reductionOperands() -> _ods_ir .. py:method:: inclusiveUpperbound() -> Optional[_ods_ir] .. py:method:: collapse() -> Optional[_ods_ir] .. py:method:: collapseDeviceType() -> Optional[_ods_ir] .. py:method:: gangOperandsArgType() -> Optional[_ods_ir] .. py:method:: gangOperandsSegments() -> Optional[_ods_ir] .. py:method:: gangOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: workerNumOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: vectorOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: seq() -> Optional[_ods_ir] .. py:method:: independent() -> Optional[_ods_ir] .. py:method:: auto_() -> Optional[_ods_ir] .. py:method:: gang() -> Optional[_ods_ir] .. py:method:: worker() -> Optional[_ods_ir] .. py:method:: vector() -> Optional[_ods_ir] .. py:method:: tileOperandsSegments() -> Optional[_ods_ir] .. py:method:: tileOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: combined() -> Optional[_ods_ir] .. py:method:: unstructured() -> bool .. py:method:: results_() -> _ods_ir .. py:method:: region() -> _ods_ir .. py:class:: LoopOpAdaptor(operands: list[Value], attributes: OpAttributeMap) LoopOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.loop' .. py:method:: lowerbound() -> _ods_ir .. py:method:: upperbound() -> _ods_ir .. py:method:: step() -> _ods_ir .. py:method:: gangOperands() -> _ods_ir .. py:method:: workerNumOperands() -> _ods_ir .. py:method:: vectorOperands() -> _ods_ir .. py:method:: tileOperands() -> _ods_ir .. py:method:: cacheOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: reductionOperands() -> _ods_ir .. py:method:: inclusiveUpperbound() -> Optional[_ods_ir] .. py:method:: collapse() -> Optional[_ods_ir] .. py:method:: collapseDeviceType() -> Optional[_ods_ir] .. py:method:: gangOperandsArgType() -> Optional[_ods_ir] .. py:method:: gangOperandsSegments() -> Optional[_ods_ir] .. py:method:: gangOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: workerNumOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: vectorOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: seq() -> Optional[_ods_ir] .. py:method:: independent() -> Optional[_ods_ir] .. py:method:: auto_() -> Optional[_ods_ir] .. py:method:: gang() -> Optional[_ods_ir] .. py:method:: worker() -> Optional[_ods_ir] .. py:method:: vector() -> Optional[_ods_ir] .. py:method:: tileOperandsSegments() -> Optional[_ods_ir] .. py:method:: tileOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: combined() -> Optional[_ods_ir] .. py:method:: unstructured() -> bool .. py:function:: loop(results_: Sequence[_ods_ir], lowerbound: Sequence[_ods_ir], upperbound: Sequence[_ods_ir], step: Sequence[_ods_ir], gang_operands: Sequence[_ods_ir], worker_num_operands: Sequence[_ods_ir], vector_operands: Sequence[_ods_ir], tile_operands: Sequence[_ods_ir], cache_operands: Sequence[_ods_ir], private_operands: Sequence[_ods_ir], firstprivate_operands: Sequence[_ods_ir], reduction_operands: Sequence[_ods_ir], *, inclusive_upperbound: Optional[Union[Sequence[bool], _ods_ir]] = None, collapse: Optional[Union[Sequence[int], _ods_ir]] = None, collapse_device_type: Optional[Union[Any, _ods_ir]] = None, gang_operands_arg_type: Optional[Union[Any, _ods_ir]] = None, gang_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, gang_operands_device_type: Optional[Union[Any, _ods_ir]] = None, worker_num_operands_device_type: Optional[Union[Any, _ods_ir]] = None, vector_operands_device_type: Optional[Union[Any, _ods_ir]] = None, seq: Optional[Union[Any, _ods_ir]] = None, independent: Optional[Union[Any, _ods_ir]] = None, auto_: Optional[Union[Any, _ods_ir]] = None, gang: Optional[Union[Any, _ods_ir]] = None, worker: Optional[Union[Any, _ods_ir]] = None, vector: Optional[Union[Any, _ods_ir]] = None, tile_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, tile_operands_device_type: Optional[Union[Any, _ods_ir]] = None, combined: Optional[Union[Any, _ods_ir]] = None, unstructured: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> Union[_ods_ir, _ods_ir, LoopOp] .. py:class:: MapInfoOp(var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, desc: Optional[_ods_ir] = None, descKind: Optional[Union[Any, _ods_ir]] = None, size: Optional[_ods_ir] = None, mapFlags: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, elementSize: Optional[Union[int, _ods_ir]] = None, exitLoc: Optional[Union[Any, _ods_ir]] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Captures the runtime allocation or mapping contract for one object, including attach points, descriptor facts, bounds, byte size, and offload map-type flags. This operation does not itself allocate or map the object. ``accVar`` is a result token of the same type as ``var``; a construct consumes the token to issue the corresponding runtime operation. Data clauses use the token as a data operand of a data, compute, declare, or update construct. When ``var`` is an ``acc.privatize`` result, the token is used as the kernel argument and its ``private`` and parallel-level map flags request the runtime allocation with the required replication. * ``var``: Mapped variable (pointer-like or mappable, same as data-entry ``var``). * ``varType``: Type of the mapped object (same role as data-entry ``varType``). * ``varPtrPtr``: Optional attach point (host address of the pointer slot). * ``desc``: Optional base descriptor when it differs from ``var``. Omitted when ``var`` itself is that descriptor; consumers then use ``var`` whenever ``descKind`` is not ``none``. * ``descKind``: Which descriptor layouts describe the mapped object, as a bitfield so that a nested descriptor can name each level. ``none`` means the object is described by its address and ``size`` alone. * ``mapFlags``: Offload map-type flags (``to``/``from``/``ptr_and_obj``/``private``/ ...), combining enter and exit clause effects or describing a privatized allocation. * ``elementSize``: Optional byte size of one element of the mapped object. ``bounds`` and descriptor extents count elements rather than bytes, so this is what converts them into byte strides. It is stated explicitly because ``varType`` does not always determine it: a character or derived-type element takes its length from the descriptor. * ``size``: Optional total mapped byte size. A constant ``0`` means the size is carried by ``bounds`` or by a descriptor instead of being stated here, a constant ``-1`` means it is not known at compile time, and a non-constant value supplies the size at run time (e.g. loaded from a type descriptor). * ``exitLoc``: Optional source location of the exit effects, which for a structured construct is its end directive rather than the location of this operation. Reported by the runtime for the region end. * ``accVar``: Result token. Its use by a construct causes that construct to issue the runtime mapping or allocation described here. .. py:attribute:: OPERATION_NAME :value: 'acc.map_info' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: desc() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: size() -> Optional[_ods_ir] .. py:method:: varType() -> _ods_ir .. py:method:: descKind() -> _ods_ir .. py:method:: mapFlags() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: elementSize() -> Optional[_ods_ir] .. py:method:: exitLoc() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: MapInfoOpAdaptor(operands: list[Value], attributes: OpAttributeMap) MapInfoOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.map_info' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: desc() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: size() -> Optional[_ods_ir] .. py:method:: varType() -> _ods_ir .. py:method:: descKind() -> _ods_ir .. py:method:: mapFlags() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: elementSize() -> Optional[_ods_ir] .. py:method:: exitLoc() -> Optional[_ods_ir] .. py:function:: map_info(var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, desc: Optional[_ods_ir] = None, desc_kind: Optional[Union[Any, _ods_ir]] = None, size: Optional[_ods_ir] = None, map_flags: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, element_size: Optional[Union[int, _ods_ir]] = None, exit_loc: Optional[Union[Any, _ods_ir]] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: NoCreateOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.nocreate' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: NoCreateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) NoCreateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.nocreate' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: nocreate(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: OnDeviceOp(deviceType: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Represents a call to the OpenACC ``acc_on_device`` runtime function. Returns whether the current thread is executing on the given device type. Example: .. code:: mlir %host = arith.constant 1 : i32 %on_host = acc.on_device %host : i32 -> i1 .. py:attribute:: OPERATION_NAME :value: 'acc.on_device' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: deviceType() -> _ods_ir .. py:method:: result() -> _ods_ir[_ods_ir] Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: OnDeviceOpAdaptor(operands: list[Value], attributes: OpAttributeMap) OnDeviceOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.on_device' .. py:method:: deviceType() -> _ods_ir .. py:function:: on_device(device_type: _ods_ir, *, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: ParWidthOp(par_dim: Union[Any, _ods_ir], *, launchArg: Optional[_ods_ir[_ods_ir]] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The ``acc.par_width`` operation specifies the parallel width for a given GPU parallel dimension. It is used as an input to ``acc.compute_region`` to define the launch configuration. The optional ``launchArg`` operand provides a known width value. When absent, the width is unknown and must be determined later (either at compile time by analysis or at runtime). Examples: .. code:: mlir // Known width from SSA value %w1 = acc.par_width %vector_len par_dim(#acc.par_dim) // Unknown width (to be computed later) %w2 = acc.par_width par_dim(#acc.par_dim) .. py:attribute:: OPERATION_NAME :value: 'acc.par_width' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: launchArg() -> Optional[_ods_ir[_ods_ir]] .. py:method:: par_dim() -> _ods_ir .. py:method:: output() -> _ods_ir[_ods_ir] .. py:class:: ParWidthOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ParWidthOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.par_width' .. py:method:: launchArg() -> Optional[_ods_ir[_ods_ir]] .. py:method:: par_dim() -> _ods_ir .. py:function:: par_width(par_dim: Union[Any, _ods_ir], *, launch_arg: Optional[_ods_ir[_ods_ir]] = None, results: Optional[Sequence[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir[_ods_ir] .. py:class:: ParallelOp(asyncOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], numGangs: Sequence[_ods_ir], numWorkers: Sequence[_ods_ir], vectorLength: Sequence[_ods_ir], reductionOperands: Sequence[_ods_ir], privateOperands: Sequence[_ods_ir], firstprivateOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, waitOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, waitOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, hasWaitDevnum: Optional[Union[Sequence[bool], _ods_ir]] = None, waitOnly: Optional[Union[Any, _ods_ir]] = None, numGangsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, numGangsDeviceType: Optional[Union[Any, _ods_ir]] = None, numWorkersDeviceType: Optional[Union[Any, _ods_ir]] = None, vectorLengthDeviceType: Optional[Union[Any, _ods_ir]] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, selfCond: Optional[_ods_ir[_ods_ir]] = None, selfAttr: Optional[bool] = None, defaultAttr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.parallel' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: numGangs() -> _ods_ir .. py:method:: numWorkers() -> _ods_ir .. py:method:: vectorLength() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: numGangsSegments() -> Optional[_ods_ir] .. py:method:: numGangsDeviceType() -> Optional[_ods_ir] .. py:method:: numWorkersDeviceType() -> Optional[_ods_ir] .. py:method:: vectorLengthDeviceType() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:method:: region() -> _ods_ir .. py:class:: ParallelOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ParallelOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.parallel' .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: numGangs() -> _ods_ir .. py:method:: numWorkers() -> _ods_ir .. py:method:: vectorLength() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: numGangsSegments() -> Optional[_ods_ir] .. py:method:: numGangsDeviceType() -> Optional[_ods_ir] .. py:method:: numWorkersDeviceType() -> Optional[_ods_ir] .. py:method:: vectorLengthDeviceType() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:function:: parallel(async_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], num_gangs: Sequence[_ods_ir], num_workers: Sequence[_ods_ir], vector_length: Sequence[_ods_ir], reduction_operands: Sequence[_ods_ir], private_operands: Sequence[_ods_ir], firstprivate_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, wait_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, wait_operands_device_type: Optional[Union[Any, _ods_ir]] = None, has_wait_devnum: Optional[Union[Sequence[bool], _ods_ir]] = None, wait_only: Optional[Union[Any, _ods_ir]] = None, num_gangs_segments: Optional[Union[Sequence[int], _ods_ir]] = None, num_gangs_device_type: Optional[Union[Any, _ods_ir]] = None, num_workers_device_type: Optional[Union[Any, _ods_ir]] = None, vector_length_device_type: Optional[Union[Any, _ods_ir]] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, self_cond: Optional[_ods_ir[_ods_ir]] = None, self_attr: Optional[bool] = None, default_attr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ParallelOp .. py:class:: PredicateRegionOp(*, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Groups statements within an ``acc.compute_region`` that sit at an intermediate point in a loop nest (outside a partitioned loop body, between or around nested loops). This grouping marks code whose execution scope differs from that of surrounding partitioned loops, so predication and synchronization can be applied correctly during lowering. Corresponds to OpenACC single or redundant execution at nest transitions. Example: .. code:: mlir // !$acc parallel num_gangs(NG) vector_length(VL) // !$acc loop gang // !$acc atomic update // !$acc loop vector // !$acc atomic update %w_gang = acc.par_width %cNG {par_dim = #acc.par_dim} %w_vector = acc.par_width %cVL {par_dim = #acc.par_dim} acc.compute_region launch(%ng = %w_gang, %vl = %w_vector) ins(%arg_c1 = %c1, %arg_c2 = %c2) : (memref, memref) { scf.parallel (%i) = (%c1) to (%cN) step (%c1) { acc.predicate_region { acc.atomic.update %arg_c1 : memref { ^bb0(%old: i32): %one = arith.constant 1 : i32 %sum = arith.addi %old, %one : i32 acc.yield %sum : i32 } } scf.parallel (%j) = (%c1) to (%cN) step (%c1) { acc.atomic.update %arg_c2 : memref { ^bb0(%old: i32): %one = arith.constant 1 : i32 %sum = arith.addi %old, %one : i32 acc.yield %sum : i32 } scf.reduce } {acc.par_dims = #acc} scf.reduce } {acc.par_dims = #acc} acc.yield } <{origin = "acc.parallel"}> .. py:attribute:: OPERATION_NAME :value: 'acc.predicate_region' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: region() -> _ods_ir .. py:class:: PredicateRegionOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PredicateRegionOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.predicate_region' .. py:function:: predicate_region(*, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> PredicateRegionOp .. py:class:: PresentOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.present' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: PresentOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PresentOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.present' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: present(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: PrivateLocalOp(output: _ods_ir, privatized: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Given a value of type ``acc.private_type``, materializes the underlying storage (typically a ``memref``) for the current thread / iteration. Which slice of the privatized allocation is selected is determined by surrounding parallelism assigned in the context. The result type is usually ``T`` (often a ``memref`` that matches the logical storage type). The result may instead use a different surface type that still aliases the same underlying storage as ``T``; in that case passes that consume this operation must treat the handle type and the result type as describing the same storage layout. .. py:attribute:: OPERATION_NAME :value: 'acc.private_local' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: privatized() -> _ods_ir .. py:method:: output() -> _ods_ir .. py:class:: PrivateLocalOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PrivateLocalOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.private_local' .. py:method:: privatized() -> _ods_ir .. py:function:: private_local(output: _ods_ir, privatized: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: PrivateOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.private' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: PrivateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PrivateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.private' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: private(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: PrivateRecipeOp(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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 first and may yield additional values that are used only for destruction. #. 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 any additional destruction values yielded by the init region, followed by 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: .. code:: mlir 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 // ... 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>) { } .. py:attribute:: OPERATION_NAME :value: 'acc.private.recipe' .. py:attribute:: _ODS_REGIONS :value: (2, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:method:: initRegion() -> _ods_ir .. py:method:: destroyRegion() -> _ods_ir .. py:class:: PrivateRecipeOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PrivateRecipeOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.private.recipe' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:function:: private_recipe(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> PrivateRecipeOp .. py:class:: PrivatizeOp(result: _ods_ir, dynamicSizes: Sequence[_ods_ir[_ods_ir]], *, par_dims: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Introduces a privatization handle for storage that varies across the active parallel dimensions (for example OpenACC ``private`` / ``firstprivate`` after recipe materialization). The handle type is ``acc.private_type`` where ``T`` is the logical storage type (commonly a ``memref``). Optional ``index`` operands supply dynamic sizes when the privatized shape depends on SSA values (for example ``memref`` row lengths). The optional ``par_dims`` attribute records which GPU parallel dimensions participate in the privatization. .. py:attribute:: OPERATION_NAME :value: 'acc.privatize' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: dynamicSizes() -> _ods_ir[_ods_ir] .. py:method:: par_dims() -> Optional[_ods_ir] .. py:method:: result() -> _ods_ir Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: PrivatizeOpAdaptor(operands: list[Value], attributes: OpAttributeMap) PrivatizeOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.privatize' .. py:method:: dynamicSizes() -> _ods_ir[_ods_ir] .. py:method:: par_dims() -> Optional[_ods_ir] .. py:function:: privatize(result: _ods_ir, dynamic_sizes: Sequence[_ods_ir[_ods_ir]], *, par_dims: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: ReductionAccumulateArrayOp(memref: _ods_ir, bounds: _ods_ir, reductionOperator: Union[Any, _ods_ir], par_dims: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Accumulates elements of an array across the specified parallel dimension given an acc.bounds op. .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_accumulate_array' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: memref() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> _ods_ir .. py:class:: ReductionAccumulateArrayOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionAccumulateArrayOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_accumulate_array' .. py:method:: memref() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> _ods_ir .. py:function:: reduction_accumulate_array(memref: _ods_ir, bounds: _ods_ir, reduction_operator: Union[Any, _ods_ir], par_dims: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ReductionAccumulateArrayOp .. py:class:: ReductionAccumulateOp(value: _ods_ir, memref: _ods_ir, reductionOperator: Union[Any, _ods_ir], par_dims: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Accumulates a scalar SSA value into a pointer-like reduction variable. Example: .. code:: mlir %private = memref.alloca() {acc.par_dims = #acc} : memref memref.store %c0, %private[] : memref %partial = scf.parallel (%iv) = (%c0) to (%cN) step (%c1) init (%c0) -> i32 { %v = memref.load %data[%iv] : memref scf.reduce(%v : i32) { ^bb0(%lhs: i32, %rhs: i32): %sum = arith.addi %lhs, %rhs : i32 scf.reduce.return %sum : i32 } } {acc.par_dims = #acc} acc.reduction_accumulate %partial to %private par_dims(#acc) : i32 -> memref acc.reduction_combine %private into %shared par_dims(#acc) : memref .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_accumulate' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: value() -> _ods_ir .. py:method:: memref() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> _ods_ir .. py:class:: ReductionAccumulateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionAccumulateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_accumulate' .. py:method:: value() -> _ods_ir .. py:method:: memref() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> _ods_ir .. py:function:: reduction_accumulate(value: _ods_ir, memref: _ods_ir, reduction_operator: Union[Any, _ods_ir], par_dims: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ReductionAccumulateOp .. py:class:: ReductionCombineOp(destMemref: _ods_ir, srcMemref: _ods_ir, reductionOperator: Union[Any, _ods_ir], *, par_dims: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` This operation is a composite to do a typical update of a reduction variable. The intention of this operator is to facilitate codegen decisions (such as generate an atomic update). E.g. .. code:: acc.reduction_combine %src into %dest par_dims(#acc) : memref Might lower to something similar to .. code:: %loadSrc = memref.load %src[] : memref %loadDest = memref.load %dest[] : memref %combine = arith.addi %loadSrc, %loadDest : i32 memref.store %combine, %dest[] : memref The ``destMemref`` operand is a "pointer" to the original reduction variable (typically shared). The ``srcMemref`` operand is a "pointer" to the partial sum of the reduction (typically private). The ``kind`` is the OpenACC reduction operator that determines how to accumulate the two values. .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_combine' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: destMemref() -> _ods_ir .. py:method:: srcMemref() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> Optional[_ods_ir] .. py:class:: ReductionCombineOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionCombineOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_combine' .. py:method:: destMemref() -> _ods_ir .. py:method:: srcMemref() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: par_dims() -> Optional[_ods_ir] .. py:function:: reduction_combine(dest_memref: _ods_ir, src_memref: _ods_ir, reduction_operator: Union[Any, _ods_ir], *, par_dims: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ReductionCombineOp .. py:class:: ReductionCombineRegionOp(destVar: _ods_ir, srcVar: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` This operation provides materialized reduction combine code from an OpenACC reduction recipe. The region takes the partially reduced value(s) from the private reduction variable and combines them with the current value(s) in the original/shared reduction variable. The region is terminated by ``acc.yield`` with no operands. The ``destVar`` operand is the original/shared reduction variable. The ``srcVar`` operand is typically the result of acc.reduction_init. .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_combine_region' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: destVar() -> _ods_ir .. py:method:: srcVar() -> _ods_ir .. py:method:: region() -> _ods_ir .. py:class:: ReductionCombineRegionOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionCombineRegionOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_combine_region' .. py:method:: destVar() -> _ods_ir .. py:method:: srcVar() -> _ods_ir .. py:function:: reduction_combine_region(dest_var: _ods_ir, src_var: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ReductionCombineRegionOp .. py:class:: ReductionInitOp(result: _ods_ir, var: _ods_ir, bounds: Sequence[_ods_ir], reductionOperator: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` This operation provides materialized allocation and initialization for a private reduction variable from an OpenACC reduction recipe. The region contains the recipe's init code and must yield a single value (the private reduction storage) via ``acc.yield``. The ``var`` operand is the original/shared reduction variable. The ``reduction_operator`` specifies the reduction kind (e.g. add, mul). The optional ``bounds`` operands describe the element range of the reduction variable (as ``acc.bounds`` ops) when it refers to an array section. They carry the section's bound information. .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_init' .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: result() -> _ods_ir Shortcut to get an op result if it has only one (throws an error otherwise). .. py:method:: region() -> _ods_ir .. py:class:: ReductionInitOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionInitOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction_init' .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:function:: reduction_init(result: _ods_ir, var: _ods_ir, bounds: Sequence[_ods_ir], reduction_operator: Union[Any, _ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: ReductionOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.reduction' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: ReductionOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: reduction(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: ReductionRecipeOp(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], reductionOperator: Union[Any, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir acc.reduction.recipe @reduction_add_memref : memref<10x20xf32> reduction_operator 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 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 %cst = arith.constant 0.0 : f32 linalg.fill ins(%cst : f32) outs(%slice_alloc : memref) // ... 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> %rhs_slice = memref.subview %rhs[%lb_outer, %lb_inner][%extent_outer, %extent_inner][1, 1] : memref<10x20xf32> to memref> // Combine only the slice portions linalg.add ins(%lhs_slice, %rhs_slice : memref>, memref>) outs(%lhs_slice : memref>) 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. | -------------------------------------------------| .. py:attribute:: OPERATION_NAME :value: 'acc.reduction.recipe' .. py:attribute:: _ODS_REGIONS :value: (3, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:method:: initRegion() -> _ods_ir .. py:method:: combinerRegion() -> _ods_ir .. py:method:: destroyRegion() -> _ods_ir .. py:class:: ReductionRecipeOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ReductionRecipeOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.reduction.recipe' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: type_() -> _ods_ir .. py:method:: reductionOperator() -> _ods_ir .. py:function:: reduction_recipe(sym_name: Union[str, _ods_ir], type_: Union[_ods_ir, _ods_ir], reduction_operator: Union[Any, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ReductionRecipeOp .. py:class:: RoutineOp(sym_name: Union[str, _ods_ir], func_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, bindIdName: Optional[Union[Any, _ods_ir]] = None, bindStrName: Optional[Union[Sequence[str], _ods_ir]] = None, bindIdNameDeviceType: Optional[Union[Any, _ods_ir]] = None, bindStrNameDeviceType: Optional[Union[Any, _ods_ir]] = None, worker: Optional[Union[Any, _ods_ir]] = None, vector: Optional[Union[Any, _ods_ir]] = None, seq: Optional[Union[Any, _ods_ir]] = None, nohost: Optional[bool] = None, implicit: Optional[bool] = None, gang: Optional[Union[Any, _ods_ir]] = None, gangDim: Optional[Union[Sequence[int], _ods_ir]] = None, gangDimDeviceType: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.routine' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: func_name() -> _ods_ir .. py:method:: bindIdName() -> Optional[_ods_ir] .. py:method:: bindStrName() -> Optional[_ods_ir] .. py:method:: bindIdNameDeviceType() -> Optional[_ods_ir] .. py:method:: bindStrNameDeviceType() -> Optional[_ods_ir] .. py:method:: worker() -> Optional[_ods_ir] .. py:method:: vector() -> Optional[_ods_ir] .. py:method:: seq() -> Optional[_ods_ir] .. py:method:: nohost() -> bool .. py:method:: implicit() -> bool .. py:method:: gang() -> Optional[_ods_ir] .. py:method:: gangDim() -> Optional[_ods_ir] .. py:method:: gangDimDeviceType() -> Optional[_ods_ir] .. py:class:: RoutineOpAdaptor(operands: list[Value], attributes: OpAttributeMap) RoutineOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.routine' .. py:method:: sym_name() -> _ods_ir .. py:method:: sym_visibility() -> Optional[_ods_ir] .. py:method:: func_name() -> _ods_ir .. py:method:: bindIdName() -> Optional[_ods_ir] .. py:method:: bindStrName() -> Optional[_ods_ir] .. py:method:: bindIdNameDeviceType() -> Optional[_ods_ir] .. py:method:: bindStrNameDeviceType() -> Optional[_ods_ir] .. py:method:: worker() -> Optional[_ods_ir] .. py:method:: vector() -> Optional[_ods_ir] .. py:method:: seq() -> Optional[_ods_ir] .. py:method:: nohost() -> bool .. py:method:: implicit() -> bool .. py:method:: gang() -> Optional[_ods_ir] .. py:method:: gangDim() -> Optional[_ods_ir] .. py:method:: gangDimDeviceType() -> Optional[_ods_ir] .. py:function:: routine(sym_name: Union[str, _ods_ir], func_name: Union[str, _ods_ir], *, sym_visibility: Optional[Union[str, _ods_ir]] = None, bind_id_name: Optional[Union[Any, _ods_ir]] = None, bind_str_name: Optional[Union[Sequence[str], _ods_ir]] = None, bind_id_name_device_type: Optional[Union[Any, _ods_ir]] = None, bind_str_name_device_type: Optional[Union[Any, _ods_ir]] = None, worker: Optional[Union[Any, _ods_ir]] = None, vector: Optional[Union[Any, _ods_ir]] = None, seq: Optional[Union[Any, _ods_ir]] = None, nohost: Optional[bool] = None, implicit: Optional[bool] = None, gang: Optional[Union[Any, _ods_ir]] = None, gang_dim: Optional[Union[Sequence[int], _ods_ir]] = None, gang_dim_device_type: Optional[Union[Any, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> RoutineOp .. py:class:: SerialOp(asyncOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], reductionOperands: Sequence[_ods_ir], privateOperands: Sequence[_ods_ir], firstprivateOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, waitOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, waitOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, hasWaitDevnum: Optional[Union[Sequence[bool], _ods_ir]] = None, waitOnly: Optional[Union[Any, _ods_ir]] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, selfCond: Optional[_ods_ir[_ods_ir]] = None, selfAttr: Optional[bool] = None, defaultAttr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.serial' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (1, True) .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:method:: region() -> _ods_ir .. py:class:: SerialOpAdaptor(operands: list[Value], attributes: OpAttributeMap) SerialOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.serial' .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: selfCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: reductionOperands() -> _ods_ir .. py:method:: privateOperands() -> _ods_ir .. py:method:: firstprivateOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: selfAttr() -> bool .. py:method:: defaultAttr() -> Optional[_ods_ir] .. py:method:: combined() -> bool .. py:function:: serial(async_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], reduction_operands: Sequence[_ods_ir], private_operands: Sequence[_ods_ir], firstprivate_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, wait_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, wait_operands_device_type: Optional[Union[Any, _ods_ir]] = None, has_wait_devnum: Optional[Union[Sequence[bool], _ods_ir]] = None, wait_only: Optional[Union[Any, _ods_ir]] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, self_cond: Optional[_ods_ir[_ods_ir]] = None, self_attr: Optional[bool] = None, default_attr: Optional[Union[Any, _ods_ir]] = None, combined: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> SerialOp .. py:class:: SetOp(*, device_type: Optional[Union[Any, _ods_ir]] = None, defaultAsync: Optional[_ods_ir] = None, deviceNum: Optional[_ods_ir] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.set" operation represents the OpenACC set directive. Example: .. code:: mlir acc.set device_num(%dev1 : i32) .. py:attribute:: OPERATION_NAME :value: 'acc.set' .. py:attribute:: _ODS_OPERAND_SEGMENTS :value: [0, 0, 0] .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: defaultAsync() -> Optional[_ods_ir] .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_type() -> Optional[_ods_ir] .. py:class:: SetOpAdaptor(operands: list[Value], attributes: OpAttributeMap) SetOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.set' .. py:method:: defaultAsync() -> Optional[_ods_ir] .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_type() -> Optional[_ods_ir] .. py:function:: set(*, device_type: Optional[Union[Any, _ods_ir]] = None, default_async: Optional[_ods_ir] = None, device_num: Optional[_ods_ir] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> SetOp .. py:class:: ShutdownOp(*, device_types: Optional[Union[Any, _ods_ir]] = None, deviceNum: Optional[_ods_ir] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.shutdown" operation represents the OpenACC shutdown executable directive. Example: .. code:: mlir acc.shutdown acc.shutdown device_num(%dev1 : i32) .. py:attribute:: OPERATION_NAME :value: 'acc.shutdown' .. py:attribute:: _ODS_OPERAND_SEGMENTS :value: [0, 0] .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_types() -> Optional[_ods_ir] .. py:class:: ShutdownOpAdaptor(operands: list[Value], attributes: OpAttributeMap) ShutdownOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.shutdown' .. py:method:: deviceNum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: device_types() -> Optional[_ods_ir] .. py:function:: shutdown(*, device_types: Optional[Union[Any, _ods_ir]] = None, device_num: Optional[_ods_ir] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> ShutdownOp .. py:class:: TerminatorOp(*, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.terminator' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:class:: TerminatorOpAdaptor(operands: list[Value], attributes: OpAttributeMap) TerminatorOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.terminator' .. py:function:: terminator(*, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> TerminatorOp .. py:class:: UnwrapPrivateOp(result: _ods_ir, handle: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` Converts a privatization handle (``acc.private_type``) to a pointer-like view of the underlying private storage. .. py:attribute:: OPERATION_NAME :value: 'acc.unwrap_private' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: handle() -> _ods_ir .. py:method:: result() -> _ods_ir Shortcut to get an op result if it has only one (throws an error otherwise). .. py:class:: UnwrapPrivateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) UnwrapPrivateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.unwrap_private' .. py:method:: handle() -> _ods_ir .. py:function:: unwrap_private(result: _ods_ir, handle: _ods_ir, *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: UpdateDeviceOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.update_device' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: UpdateDeviceOpAdaptor(operands: list[Value], attributes: OpAttributeMap) UpdateDeviceOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.update_device' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: update_device(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: UpdateHostOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.update_host' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: accVar() -> _ods_ir .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:class:: UpdateHostOpAdaptor(operands: list[Value], attributes: OpAttributeMap) UpdateHostOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.update_host' .. py:method:: accVar() -> _ods_ir .. py:method:: var() -> _ods_ir .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:function:: update_host(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> UpdateHostOp .. py:class:: UpdateOp(asyncOperands: Sequence[_ods_ir], waitOperands: Sequence[_ods_ir], dataClauseOperands: Sequence[_ods_ir], *, ifCond: Optional[_ods_ir[_ods_ir]] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, waitOperandsSegments: Optional[Union[Sequence[int], _ods_ir]] = None, waitOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, hasWaitDevnum: Optional[Union[Sequence[bool], _ods_ir]] = None, waitOnly: Optional[Union[Any, _ods_ir]] = None, ifPresent: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.update' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: ifPresent() -> bool .. py:class:: UpdateOpAdaptor(operands: list[Value], attributes: OpAttributeMap) UpdateOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.update' .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: asyncOperands() -> _ods_ir .. py:method:: waitOperands() -> _ods_ir .. py:method:: dataClauseOperands() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: waitOperandsSegments() -> Optional[_ods_ir] .. py:method:: waitOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: hasWaitDevnum() -> Optional[_ods_ir] .. py:method:: waitOnly() -> Optional[_ods_ir] .. py:method:: ifPresent() -> bool .. py:function:: update(async_operands: Sequence[_ods_ir], wait_operands: Sequence[_ods_ir], data_clause_operands: Sequence[_ods_ir], *, if_cond: Optional[_ods_ir[_ods_ir]] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, wait_operands_segments: Optional[Union[Sequence[int], _ods_ir]] = None, wait_operands_device_type: Optional[Union[Any, _ods_ir]] = None, has_wait_devnum: Optional[Union[Sequence[bool], _ods_ir]] = None, wait_only: Optional[Union[Any, _ods_ir]] = None, if_present: Optional[bool] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> UpdateOp .. py:class:: UseDeviceOp(accVar: _ods_ir, var: _ods_ir, varType: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], asyncOperands: Sequence[_ods_ir], *, varPtrPtr: Optional[_ods_ir] = None, asyncOperandsDeviceType: Optional[Union[Any, _ods_ir]] = None, asyncOnly: Optional[Union[Any, _ods_ir]] = None, dataClause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.use_device' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] Returns the fully qualified name of the operation. .. py:method:: recipe() -> Optional[_ods_ir] .. py:method:: accVar() -> _ods_ir .. py:class:: UseDeviceOpAdaptor(operands: list[Value], attributes: OpAttributeMap) UseDeviceOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.use_device' .. py:method:: var() -> _ods_ir .. py:method:: varPtrPtr() -> Optional[_ods_ir] .. py:method:: bounds() -> _ods_ir .. py:method:: asyncOperands() -> _ods_ir .. py:method:: varType() -> _ods_ir .. py:method:: asyncOperandsDeviceType() -> Optional[_ods_ir] .. py:method:: asyncOnly() -> Optional[_ods_ir] .. py:method:: dataClause() -> _ods_ir .. py:method:: structured() -> _ods_ir .. py:method:: implicit() -> _ods_ir .. py:method:: modifiers() -> _ods_ir .. py:method:: name() -> Optional[_ods_ir] .. py:method:: recipe() -> Optional[_ods_ir] .. py:function:: use_device(acc_var: _ods_ir, var: _ods_ir, var_type: Union[_ods_ir, _ods_ir], bounds: Sequence[_ods_ir], async_operands: Sequence[_ods_ir], *, var_ptr_ptr: Optional[_ods_ir] = None, async_operands_device_type: Optional[Union[Any, _ods_ir]] = None, async_only: Optional[Union[Any, _ods_ir]] = None, data_clause: Optional[Union[Any, _ods_ir]] = None, structured: Optional[Union[bool, _ods_ir]] = None, implicit: Optional[Union[bool, _ods_ir]] = None, modifiers: Optional[Union[Any, _ods_ir]] = None, name: Optional[Union[str, _ods_ir]] = None, recipe: Optional[Union[str, _ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> _ods_ir .. py:class:: WaitOp(waitOperands: Sequence[_ods_ir], *, asyncOperand: Optional[_ods_ir] = None, waitDevnum: Optional[_ods_ir] = None, async_: Optional[bool] = None, ifCond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_ods_ir` The "acc.wait" operation represents the OpenACC wait executable directive. Example: .. code:: mlir 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. .. py:attribute:: OPERATION_NAME :value: 'acc.wait' .. py:attribute:: _ODS_OPERAND_SEGMENTS .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: waitOperands() -> _ods_ir .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: async_() -> bool .. py:class:: WaitOpAdaptor(operands: list[Value], attributes: OpAttributeMap) WaitOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.wait' .. py:method:: waitOperands() -> _ods_ir .. py:method:: asyncOperand() -> Optional[_ods_ir] .. py:method:: waitDevnum() -> Optional[_ods_ir] .. py:method:: ifCond() -> Optional[_ods_ir[_ods_ir]] .. py:method:: async_() -> bool .. py:function:: wait(wait_operands: Sequence[_ods_ir], *, async_operand: Optional[_ods_ir] = None, wait_devnum: Optional[_ods_ir] = None, async_: Optional[bool] = None, if_cond: Optional[_ods_ir[_ods_ir]] = None, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> WaitOp .. py:class:: YieldOp(operands_: Sequence[_ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) Bases: :py:obj:`_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. .. py:attribute:: OPERATION_NAME :value: 'acc.yield' .. py:attribute:: _ODS_REGIONS :value: (0, True) .. py:method:: operands_() -> _ods_ir .. py:class:: YieldOpAdaptor(operands: list[Value], attributes: OpAttributeMap) YieldOpAdaptor(operands: list[Value], opview: OpView) Bases: :py:obj:`_ods_ir` .. py:attribute:: OPERATION_NAME :value: 'acc.yield' .. py:method:: operands_() -> _ods_ir .. py:function:: yield_(operands_: Sequence[_ods_ir], *, loc: Optional[_ods_ir] = None, ip: Optional[_ods_ir] = None) -> YieldOp