CINXE.COM

<!doctype html><html lang=en-us><head><meta charset=utf-8><meta http-equiv=x-ua-compatible content="IE=edge"><meta name=viewport content="width=device-width,initial-scale=1,maximum-scale=1,user-scalable=no"><title>'acc' Dialect - MLIR</title><meta name=description content="Multi-Level IR Compiler Framework"><meta name=generator content="Hugo 0.119.0"><link href=https://mlir.llvm.org/index.xml rel=alternate type=application/rss+xml><link rel=canonical href=https://mlir.llvm.org/docs/Dialects/OpenACCDialect/><link rel=stylesheet href=https://mlir.llvm.org/css/theme.css><script src=https://use.fontawesome.com/releases/v5.0.6/js/all.js></script> <link rel=stylesheet href=https://mlir.llvm.org/css/chroma.min.css><script src=https://cdn.jsdelivr.net/npm/jquery@3.3.1/dist/jquery.min.js></script> <script src=https://cdn.jsdelivr.net/npm/jquery.easing@1.4.1/jquery.easing.min.js></script> <script src=https://mlir.llvm.org/js/bundle.js></script> <script type=text/javascript src="https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.1/MathJax.js?config=TeX-AMS-MML_HTMLorMML"></script> <script type=text/x-mathjax-config> MathJax.Hub.Config({ tex2jax: { inlineMath: [['$', '$'] ], displayMath: [ ['$$','$$'], ["\\[","\\]"] ] } }); </script><link rel=apple-touch-icon sizes=180x180 href="/apple-touch-icon.png?v=1"><link rel=icon type=image/png sizes=32x32 href="/favicon-32x32.png?v=1"><link rel=icon type=image/png sizes=16x16 href="/favicon-16x16.png?v=1"><link rel=manifest href="/site.webmanifest?v=1"><link rel=mask-icon href="/safari-pinned-tab.svg?v=1" color=#3775e0><link rel="shortcut icon" href="/favicon.ico?v=1"><meta name=msapplication-TileColor content="#2d89ef"><meta name=theme-color content="#ffffff"><link rel=icon href=/favicon.svg type=image/svg+xml sizes=any><style>:root{}</style></head><body><div class=container><header><h1><div><img src=https://mlir.llvm.org//mlir-logo.png width=40px align=absmiddle> MLIR</div></h1><p class=description>Multi-Level IR Compiler Framework</p></header><div class=global-menu><nav><ul><li class=parent><a href>Community<i class="fas fa-angle-right"></i></a><ul class=sub-menu><li class=child><a href=https://llvm.discourse.group/c/mlir/31>Forums</a></li><li class=child><a href=https://discord.gg/xS7Z362>Chat</a></li></ul></li><li><a href=/getting_started/Debugging/>Debugging Tips</a></li><li><a href=/getting_started/Faq/>FAQ</a></li><li class=parent><a href=https://github.com/llvm/llvm-project/tree/main/mlir>Source<i class="fas fa-angle-right"></i></a><ul class=sub-menu><li class=child><a href=/doxygen/>Doxygen</a></li><li class=child><a href=https://github.com/llvm/llvm-project/tree/main/mlir>GitHub</a></li></ul></li><li><a href="https://bugs.llvm.org/buglist.cgi?bug_status=__open__&list_id=177877&order=changeddate%20DESC%2Cpriority%2Cbug_severity&product=MLIR&query_format=specific">Bugs</a></li><li><a href=https://github.com/llvm/mlir-www/tree/main/website/static/LogoAssets>Logo Assets</a></li><li><a href=https://www.youtube.com/MLIRCompiler>Youtube Channel</a></li></ul></nav></div><div class=content-container><main><h1>'acc' Dialect</h1><p>The <code>acc</code> dialect is an MLIR dialect for representing the OpenACC programming model. OpenACC is a standardized directive-based model which is used with C, C++, and Fortran to enable programmers to expose parallelism in their code. The descriptive approach used by OpenACC allows targeting of parallel multicore and accelerator targets like GPUs by giving the compiler the freedom of how to parallelize for specific architectures. OpenACC also provides the ability to optimize the parallelism through increasingly more prescriptive clauses.</p><p>This dialect models the constructs from the <a href=https://www.openacc.org/sites/default/files/inline-images/Specification/OpenACC-3.3-final.pdf>OpenACC 3.3 specification</a></p><p>This document describes the design of the OpenACC dialect in MLIR. It lists and explains design goals and design choices along with their rationale. It also describes specifics with regards to acc dialect operations, types, and attributes.</p><p><nav id=TableOfContents><ul><li><a href=#dialect-design-goals>Dialect Design Goals</a></li><li><a href=#operation-categories>Operation Categories</a><ul><li><a href=#data-operations>Data Operations</a></li><li><a href=#types>Types</a></li><li><a href=#recipes>Recipes</a></li><li><a href=#routine>Routine</a></li><li><a href=#declare>Declare</a></li></ul></li><li><a href=#openacc-transforms-and-analyses>OpenACC Transforms and Analyses</a><ul><li><a href=#verification>Verification</a></li><li><a href=#implicit-data-attributes>Implicit Data Attributes</a></li><li><a href=#redundant-clause-elimination>Redundant Clause Elimination</a></li></ul></li><li><a href=#operations-toc>Operations TOC</a></li><li><a href=#operations>Operations</a><ul><li><a href=#accatomiccapture-accatomiccaptureop><code>acc.atomic.capture</code> (acc::AtomicCaptureOp)</a></li><li><a href=#accatomicread-accatomicreadop><code>acc.atomic.read</code> (acc::AtomicReadOp)</a></li><li><a href=#accatomicupdate-accatomicupdateop><code>acc.atomic.update</code> (acc::AtomicUpdateOp)</a></li><li><a href=#accatomicwrite-accatomicwriteop><code>acc.atomic.write</code> (acc::AtomicWriteOp)</a></li><li><a href=#accattach-accattachop><code>acc.attach</code> (acc::AttachOp)</a></li><li><a href=#accbounds-accdataboundsop><code>acc.bounds</code> (acc::DataBoundsOp)</a></li><li><a href=#acccache-acccacheop><code>acc.cache</code> (acc::CacheOp)</a></li><li><a href=#acccopyin-acccopyinop><code>acc.copyin</code> (acc::CopyinOp)</a></li><li><a href=#acccopyout-acccopyoutop><code>acc.copyout</code> (acc::CopyoutOp)</a></li><li><a href=#acccreate-acccreateop><code>acc.create</code> (acc::CreateOp)</a></li><li><a href=#accdata-accdataop><code>acc.data</code> (acc::DataOp)</a></li><li><a href=#accdeclare-accdeclareop><code>acc.declare</code> (acc::DeclareOp)</a></li><li><a href=#accdeclare_device_resident-accdeclaredeviceresidentop><code>acc.declare_device_resident</code> (acc::DeclareDeviceResidentOp)</a></li><li><a href=#accdeclare_enter-accdeclareenterop><code>acc.declare_enter</code> (acc::DeclareEnterOp)</a></li><li><a href=#accdeclare_exit-accdeclareexitop><code>acc.declare_exit</code> (acc::DeclareExitOp)</a></li><li><a href=#accdeclare_link-accdeclarelinkop><code>acc.declare_link</code> (acc::DeclareLinkOp)</a></li><li><a href=#accdelete-accdeleteop><code>acc.delete</code> (acc::DeleteOp)</a></li><li><a href=#accdetach-accdetachop><code>acc.detach</code> (acc::DetachOp)</a></li><li><a href=#accdeviceptr-accdeviceptrop><code>acc.deviceptr</code> (acc::DevicePtrOp)</a></li><li><a href=#accenter_data-accenterdataop><code>acc.enter_data</code> (acc::EnterDataOp)</a></li><li><a href=#accexit_data-accexitdataop><code>acc.exit_data</code> (acc::ExitDataOp)</a></li><li><a href=#accfirstprivate-accfirstprivateop><code>acc.firstprivate</code> (acc::FirstprivateOp)</a></li><li><a href=#accfirstprivaterecipe-accfirstprivaterecipeop><code>acc.firstprivate.recipe</code> (acc::FirstprivateRecipeOp)</a></li><li><a href=#accgetdeviceptr-accgetdeviceptrop><code>acc.getdeviceptr</code> (acc::GetDevicePtrOp)</a></li><li><a href=#accglobal_ctor-accglobalconstructorop><code>acc.global_ctor</code> (acc::GlobalConstructorOp)</a></li><li><a href=#accglobal_dtor-accglobaldestructorop><code>acc.global_dtor</code> (acc::GlobalDestructorOp)</a></li><li><a href=#acchost_data-acchostdataop><code>acc.host_data</code> (acc::HostDataOp)</a></li><li><a href=#accinit-accinitop><code>acc.init</code> (acc::InitOp)</a></li><li><a href=#acckernels-acckernelsop><code>acc.kernels</code> (acc::KernelsOp)</a></li><li><a href=#accloop-accloopop><code>acc.loop</code> (acc::LoopOp)</a></li><li><a href=#accnocreate-accnocreateop><code>acc.nocreate</code> (acc::NoCreateOp)</a></li><li><a href=#accparallel-accparallelop><code>acc.parallel</code> (acc::ParallelOp)</a></li><li><a href=#accpresent-accpresentop><code>acc.present</code> (acc::PresentOp)</a></li><li><a href=#accprivate-accprivateop><code>acc.private</code> (acc::PrivateOp)</a></li><li><a href=#accprivaterecipe-accprivaterecipeop><code>acc.private.recipe</code> (acc::PrivateRecipeOp)</a></li><li><a href=#accreduction-accreductionop><code>acc.reduction</code> (acc::ReductionOp)</a></li><li><a href=#accreductionrecipe-accreductionrecipeop><code>acc.reduction.recipe</code> (acc::ReductionRecipeOp)</a></li><li><a href=#accroutine-accroutineop><code>acc.routine</code> (acc::RoutineOp)</a></li><li><a href=#accserial-accserialop><code>acc.serial</code> (acc::SerialOp)</a></li><li><a href=#accset-accsetop><code>acc.set</code> (acc::SetOp)</a></li><li><a href=#accshutdown-accshutdownop><code>acc.shutdown</code> (acc::ShutdownOp)</a></li><li><a href=#accterminator-accterminatorop><code>acc.terminator</code> (acc::TerminatorOp)</a></li><li><a href=#accupdate-accupdateop><code>acc.update</code> (acc::UpdateOp)</a></li><li><a href=#accupdate_device-accupdatedeviceop><code>acc.update_device</code> (acc::UpdateDeviceOp)</a></li><li><a href=#accupdate_host-accupdatehostop><code>acc.update_host</code> (acc::UpdateHostOp)</a></li><li><a href=#accuse_device-accusedeviceop><code>acc.use_device</code> (acc::UseDeviceOp)</a></li><li><a href=#accwait-accwaitop><code>acc.wait</code> (acc::WaitOp)</a></li><li><a href=#accyield-accyieldop><code>acc.yield</code> (acc::YieldOp)</a></li></ul></li><li><a href=#attributes-40>Attributes</a><ul><li><a href=#declareactionattr>DeclareActionAttr</a></li><li><a href=#declareattr>DeclareAttr</a></li><li><a href=#clausedefaultvalueattr>ClauseDefaultValueAttr</a></li><li><a href=#combinedconstructstypeattr>CombinedConstructsTypeAttr</a></li><li><a href=#constructattr>ConstructAttr</a></li><li><a href=#dataclauseattr>DataClauseAttr</a></li><li><a href=#devicetypeattr>DeviceTypeAttr</a></li><li><a href=#gangargtypeattr>GangArgTypeAttr</a></li><li><a href=#reductionoperatorattr>ReductionOperatorAttr</a></li><li><a href=#routineinfoattr>RoutineInfoAttr</a></li></ul></li><li><a href=#types-1>Types</a><ul><li><a href=#databoundstype>DataBoundsType</a></li><li><a href=#declaretokentype>DeclareTokenType</a></li></ul></li><li><a href=#enums>Enums</a><ul><li><a href=#clausedefaultvalue>ClauseDefaultValue</a></li><li><a href=#combinedconstructstype>CombinedConstructsType</a></li><li><a href=#construct>Construct</a></li><li><a href=#dataclause>DataClause</a></li><li><a href=#devicetype>DeviceType</a></li><li><a href=#gangargtype>GangArgType</a></li><li><a href=#reductionoperator>ReductionOperator</a></li></ul></li></ul></nav><h2 id=dialect-design-goals>Dialect Design Goals <a class=headline-hash href=#dialect-design-goals>¶</a></h2><ul><li>Needs to have complete representation of the OpenACC language.<ul><li>A frontend requires this in order to properly generate a representation of possible <code>acc</code> pragmas in MLIR. Additionally, this dialect is expected to be further lowered when materializing its semantics. Without a complete representation, a frontend might choose a lower abstraction (such as direct runtime call) - but this would impact the ability to do analysis and optimizations on the dialect.</li></ul></li><li>Allow representation at the same semantic level as the OpenACC language while having capability to represent nuances of the source language semantics (such as Fortran descriptors) in an agnostic manner.<ul><li>Using abstractions that closely model the OpenACC language simplifies frontend implementation. It also allows for easier debugging of the IR. However, sometimes source language specific behavior is needed when materializing OpenACC. In these cases, such as privatization of C++ objects with default constructor, the frontend fills in the <code>recipe</code> along with the <code>private</code> operation which can be packaged neatly with the <code>acc</code> dialect operations.</li></ul></li><li>Be able to regenerate the semantic equivalent of the user pragmas from the dialect (including bounds, names, clauses, modifiers, etc).<ul><li>This is a strong measure of making sure that the dialect is not lossy in semantics. It also allows capability to generate appropriate and useful debug information outside of the frontend.</li></ul></li><li>Be dialect agnostic so that it can be used and coexist with other dialects including but not limited to <code>hlfir</code>, <code>fir</code>, <code>llvm</code>, <code>cir</code>.<ul><li>Directive-based models such as OpenACC are always used with a source language, so the <code>acc</code> dialect coexisting with other dialect(s) is necessary by construction. Through proper abstractions, neither the <code>acc</code> dialect nor the source language dialect should have dependencies on each other; where needed, interfaces should be used to ensure <code>acc</code> dialect can verify expected properties.</li></ul></li><li>The dialect must allow dataflow to be modeled accurately and performantly using MLIR’s existing facilities.<ul><li>Appropriate dataflow modeling is important for analyses and IR reasoning - even something as simple as walking the uses. Therefore operations, like data operations, are expected to generate results which can be used in modeling behavior. For example, consider an <code>acc copyin</code> clause. After the <code>acc.copyin</code> operation, a pointer which lives on devices should be distinguishable from one that lives in host memory.</li></ul></li><li>Be friendly to MLIR optimization passes by implementing common interfaces.<ul><li>Interfaces, such as <code>MemoryEffects</code>, are the key way MLIR transformations and analyses are designed to interact with the IR. In order for the operations in the <code>acc</code> dialect to be optimizable (either directly or even indirectly by not blocking optimizations of nested IR), implementing relevant common interfaces is needed.</li></ul></li></ul><p>The design philosophy of the acc dialect is one where the design goals are adhered to. Current and planned operations, attributes, types must adhere to the design goals.</p><h2 id=operation-categories>Operation Categories <a class=headline-hash href=#operation-categories>¶</a></h2><p>The OpenACC dialect includes both high-level operations (which retain the same semantic meaning as their OpenACC language equivalent), intermediate-level operations (which are used to decompose clauses from constructs), and low-level operations (to encode specifics associated with source language in a generic way).</p><p>The high-level operations list contains the following OpenACC language constructs and their corresponding operations:</p><ul><li><code>acc parallel</code> → <code>acc.parallel</code></li><li><code>acc kernels</code> → <code>acc.kernels</code></li><li><code>acc serial</code> → <code>acc.serial</code></li><li><code>acc data</code> → <code>acc.data</code></li><li><code>acc loop</code> → <code>acc.loop</code></li><li><code>acc enter data</code> → <code>acc.enter_data</code></li><li><code>acc exit data</code> → <code>acc.exit_data</code></li><li><code>acc host_data</code> → <code>acc.host_data</code></li><li><code>acc init</code> → <code>acc.init</code></li><li><code>acc shutdown</code> → <code>acc.shutdown</code></li><li><code>acc update</code> → <code>acc.update</code></li><li><code>acc set</code> → <code>acc.set</code></li><li><code>acc wait</code> → <code>acc.wait</code></li><li><code>acc atomic read</code> → <code>acc.atomic.read</code></li><li><code>acc atomic write</code> → <code>acc.atomic.write</code></li><li><code>acc atomic update</code> → <code>acc.atomic.update</code></li><li><code>acc atomic capture</code> → <code>acc.atomic.capture</code></li></ul><p>This second group contains operations which are used to represent either decomposed constructs or clauses for more accurate modeling:</p><ul><li><code>acc routine</code> → <code>acc.routine</code> + <code>acc.routine_info</code> attribute</li><li><code>acc declare</code> → <code>acc.declare_enter</code> + <code>acc.declare_exit</code> or <code>acc.declare</code></li><li><code>acc {construct} copyin</code> → <code>acc.copyin</code> (before region) + <code>acc.delete</code> (after region)</li><li><code>acc {construct} copy</code> → <code>acc.copyin</code> (before region) + <code>acc.copyout</code> (after region)</li><li><code>acc {construct} copyout</code> → <code>acc.create</code> (before region) + <code>acc.copyout</code> (after region)</li><li><code>acc {construct} attach</code> → <code>acc.attach</code> (before region) + <code>acc.detach</code> (after region)</li><li><code>acc {construct} create</code> → <code>acc.create</code> (before region) + <code>acc.delete</code> (after region)</li><li><code>acc {construct} present</code> → <code>acc.present</code> (before region) + <code>acc.delete</code> (after region)</li><li><code>acc {construct} no_create</code> → <code>acc.nocreate</code> (before region) + <code>acc.delete</code> (after region)</li><li><code>acc {construct} deviceptr</code> → <code>acc.deviceptr</code></li><li><code>acc {construct} private</code> → <code>acc.private</code></li><li><code>acc {construct} firstprivate</code> → <code>acc.firstprivate</code></li><li><code>acc {construct} reduction</code> → <code>acc.reduction</code></li><li><code>acc cache</code> → <code>acc.cache</code></li><li><code>acc update device</code> → <code>acc.update_device</code></li><li><code>acc update host</code> → <code>acc.update_host</code></li><li><code>acc host_data use_device</code> → <code>acc.use_device</code></li><li><code>acc declare device_resident</code> → <code>acc.declare_device_resident</code></li><li><code>acc declare link</code> → <code>acc.declare_link</code></li><li><code>acc exit data delete</code> → <code>acc.delete</code> (with <code>structured</code> flag as false)</li><li><code>acc exit data detach</code> → <code>acc.detach</code> (with <code>structured</code> flag as false)</li><li><code>acc {construct} {data_clause}(var[lb:ub])</code> → <code>acc.bounds</code></li></ul><p>The low-level operations are:</p><ul><li><code>acc.private.recipe</code></li><li><code>acc.reduction.recipe</code></li><li><code>acc.firstprivate.recipe</code></li><li><code>acc.global_ctor</code></li><li><code>acc.global_dtor</code></li><li><code>acc.yield</code></li><li><code>acc.terminator</code> The low-level operations semantics and reasoning are further explained in sections below.</li></ul><h3 id=data-operations>Data Operations <a class=headline-hash href=#data-operations>¶</a></h3><h4 id=data-clause-decomposition>Data Clause Decomposition <a class=headline-hash href=#data-clause-decomposition>¶</a></h4><p>The data clauses are decomposed from their constructs for better dataflow modeling in MLIR. There are multiple reasons for this which are consistent with the dialect goals:</p><ul><li>Correctly represents dataflow. Data clauses have different effects at entry to region and at exit from region.</li><li>Friendlier to add attributes such as <code>MemoryEffects</code> to a single operation. This can better reflect semantics (like the fact that an <code>acc.copyin</code> operation only reads host memory)</li><li>Operations can be moved or optimized individually (eg <code>CSE</code>).</li><li>Easier to keep track of debug information. Line location can point to the text representing the data clause instead of the construct. Additionally, attributes can be used to keep track of variable names in clauses without having to walk the IR tree in attempt to recover the information (this makes acc dialect more agnostic with regards to what other dialect it is used with).</li><li>Clear operation ordering since all data operations are on same list.</li></ul><p>Each of the <code>acc</code> dialect data operations represents either the entry or the exit portion of the data action specification. Thus, <code>acc.copyin</code> represents the semantics defined in section <code>2.7.7 copyin clause</code> whose wording starts with <code>At entry to a region</code>. The decomposed exit operation <code>acc.delete</code> represents the second part of that section, whose wording starts with <code>At exit from the region</code>. The <code>delete</code> action may be performed after checking and updating of the relevant reference counters noted.</p><p>The <code>acc</code> data operations, even when decomposed, retain their original data clause in an operation operand <code>dataClause</code> for possibility to recover this information during debugging. For example, <code>acc copy</code>, does not translate to <code>acc.copy</code> operation, but instead to <code>acc.copyin</code> for entry and <code>acc.copyout</code> for exit. Both the decomposed operations hold a <code>dataClause</code> field that specifies this was an <code>acc copy</code>.</p><p>The link between the decomposed entry and exit operations is the ssa value produced by the entry operation. Namely, it is the <code>accPtr</code> result which is used both in the <code>dataOperands</code> of the operation used for the construct and in the <code>accPtr</code> operand of the exit operation.</p><h4 id=bounds>Bounds <a class=headline-hash href=#bounds>¶</a></h4><p>OpenACC data clauses allow the use of bounds specifiers as per <code>2.7.1 Data Specification in Data Clauses</code>. However, array dimensions for the data are not always required in the clause if the source language’s type system captures this information - the user can just specify the variable name in the data clause. So the <code>acc.bounds</code> operation is an important piece to ensure uniform representation of both explicit user set dimensions and implicit type-based dimensions. It contains several key features to allow properly encoding sizes in a manner flexible and agnostic to the source language’s dialect:</p><ul><li>Multi-dimensional arrays can be represented by using multiple ordered <code>acc.bounds</code> operations.</li><li>Bounds are required to be zero-normalized. This works well with the <code>PointerLikeType</code> requirement in data clauses - since a lowerbound of 0 means looking at data at the zero offset from pointer. This requirement also works well in ensuring the <code>acc</code> dialect is agnostic to source language dialect since it prevents ambiguity such as the case of Fortran arrays where the lower bound is not a fixed value.</li><li>If the source dialect does not encode the dimensions in the type (eg <code>!fir.array<?x?xi32></code>) but instead encodes it in some other way (such as through descriptors), then the frontend must fill in the <code>acc.bounds</code> operands with appropriate information (such as loads from descriptor). The <code>acc.bounds</code> operation also permits lossy source dialect, such as if the frontend uses aggressive pointer decay and cannot represent the dimensions in the type system (eg using <code>!llvm.ptr</code> for arrays). Both of these aspects show <code>acc.bounds</code>’ operation’s flexibility to allow the representation to be agnostic since the <code>acc</code> dialect is not expected to be able to understand how to extract dimension information from the types of the source dialect.</li><li>The OpenACC specification allows either extent or upperbound in the data clause depending on whether it is Fortran or C and C++. The <code>acc.bounds</code> operation is rich enough to accept either or both - for convenience in lowering to the dialect and for ability to precisely capture the meaning from the clause.</li><li>The stride, either in units or bytes, can be also captured in the <code>acc.bounds</code> operation. This is also an important part to be able to accept a source language’s arrays without forcing the frontend to normalize them in some way. For example, consider a case where in a parent function, a whole array is mapped to device. Then only a view of a non-1 stride is passed to child function (eg Fortran array slice with non-1 stride). A <code>copy</code> operation of this data in child should be able to avoid remapping this array. If instead the operation required normalizing the array (such as making it contiguous), then unexpected disjoint mapping of the same host data would be error-prone since it would result in multiple mappings to device.</li></ul><h4 id=counters>Counters <a class=headline-hash href=#counters>¶</a></h4><p>The data operations also maintain semantics described in the OpenACC specification related to runtime counters. More specifically, consider the specification of the entry portion of <code>acc copyin</code> in section 2.7.7:</p><pre tabindex=0><code>At entry to a region, the structured reference counter is used. On an enter data directive, the dynamic reference counter is used. - If var is present and is not a null pointer, a present increment action with the appropriate reference counter is performed. - If var is not present, a copyin action with the appropriate reference counter is performed. - If var is a pointer reference, an attach action is performed. </code></pre><p>The <code>acc.copyin</code> operation includes these semantics, including those related to attach, which is specified through the <code>varPtrPtr</code> operand. The <code>structured</code> flag on the operation is important since the <code>structured reference counter</code> should be used when the flag is true; and the <code>dynamic reference counter</code> should be used when it is false.</p><p>At exit from structured regions (<code>acc data</code>, <code>acc kernels</code>), the <code>acc copyin</code> operation is decomposed to <code>acc.delete</code> (with the <code>structured</code> flag as true). The semantics of the <code>acc.delete</code> are also consistent with the OpenACC specification noted for the exit portion of the <code>acc copyin</code> clause:</p><pre tabindex=0><code>At exit from the region: - If the structured reference counter for var is zero, no action is taken. - Otherwise, a detach action is performed if var is a pointer reference, and a present decrement action with the structured reference counter is performed if var is not a null pointer. If both structured and dynamic reference counters are zero, a delete action is performed. </code></pre><h3 id=types>Types <a class=headline-hash href=#types>¶</a></h3><p>There are a few acc dialect type categories to describe:</p><ul><li>type of acc data clause operation input <code>varPtr</code><ul><li>The type of <code>varPtr</code> must be pointer-like. This is done by attaching the <code>PointerLikeType</code> interface to the appropriate MLIR type. Although memory/storage concept is a lower level abstraction, it is useful because the OpenACC model distinguishes between host and device memory explicitly - and the mapping between the two is done through pointers. Thus, by explicitly requiring it in the dialect, the appropriate language frontend must create storage or use type that satisfies the mapping constraint.</li></ul></li><li>type of result of acc data clause operations<ul><li>The type of the acc data clause operation is exactly the same as <code>varPtr</code>. This was done intentionally instead of introducing an <code>acc.ref/ptr</code> type so that IR compatibility and the dialect’s existing strong type checking can be maintained. This is needed since the <code>acc</code> dialect must live within another dialect whose type system is unknown to it. The only constraint is that the appropriate dialect type must use the <code>PointerLikeType</code> interface.</li></ul></li><li>type of decomposed clauses<ul><li>Decomposed clauses, such as <code>acc.bounds</code> and <code>acc.declare_enter</code> produce types to allow their results to be used only in specific operations.</li></ul></li></ul><h3 id=recipes>Recipes <a class=headline-hash href=#recipes>¶</a></h3><p>Recipes are a generic way to express source language specific semantics.</p><p>There are currently two categories of recipes, but the recipe concept can be extended for any additional low-level information that needs to be captured for successful lowering of OpenACC. The two categories are:</p><ul><li>recipes used in the context of privatization associated with a construct</li><li>recipes used in the context of additional specification of data semantics</li></ul><p>The intention of the recipes is to specify how materialization of action, such as privatization, should be done when the semantics of the action needs interpreted and lowered, such as before generating LLVM dialect.</p><p>The recipes used for privatization provide a source-language independent way of specifying the creation of a local variable of that type. This means using the appropriate <code>alloca</code> instruction and being able to specify default initialization or default constructor.</p><h3 id=routine>Routine <a class=headline-hash href=#routine>¶</a></h3><p>The routine directive is used to note that a procedure should be made available for the accelerator in a way that is consistent with its modifiers, such as those that describe the parallelism. In the acc dialect, an acc routine is represented through two joint pieces - an attribute and an operation:</p><ul><li>The <code>acc.routine</code> operation is simply a specifier which notes which symbol (or string) the acc routine is needed for, along with parallelism associated. This defines a symbol that can be referenced in attribute.</li><li>The <code>acc.routine_info</code> attribute is an attribute used on the source dialect specific operation which specifies one or multiple <code>acc.routine</code> symbols. Typically, this is attached to <code>func.func</code> which either provides the declaration (in case of externals) or provides the actual body of the acc routine in the dialect that the source language was translated to.</li></ul><h3 id=declare>Declare <a class=headline-hash href=#declare>¶</a></h3><p>OpenACC <code>declare</code> is a mechanism which declares a definition of a global or a local to be accessible to accelerator with an implicit lifetime as that of the scope where it was declared in. Thus, <code>declare</code> semantics are represented through multiple operations and attributes:</p><ul><li><code>acc.declare</code> - This is a structured operation which contains an MLIR region and can be used in similar manner as acc.data to specify an implicit data region with specific procedure lifetime. This is typically used inside <code>func.func</code> after variable declarations.</li><li><code>acc.declare_enter</code> - This is an unstructured operation which is used as a decomposed form of <code>acc declare</code>. It effectively allows the entry operation to exist in a scope different than the exit operation. It can also be used along <code>acc.declare_exit</code> which consumes its token to define a scoped region without using MLIR region. This operation is also used in <code>acc.global_ctor</code>.</li><li><code>acc.declare_exit</code> - The matching equivalent of <code>acc.declare_enter</code> except that it specifies exit semantics. This operation is typically used inside a <code>func.func</code> at the exit points or with <code>acc.global_dtor</code>.</li><li><code>acc.global_ctor</code> - Lives at the same level as source dialect globals and is used to specify data actions to be done at program entry. This is used in conjunction with source dialect globals whose lifetime is not just a single procedure.</li><li><code>acc.global_dtor</code> - Defines the exit data actions that should be done at program exit. Typically used to revert the actions of <code>acc.global_ctor</code>.</li></ul><p>The attributes:</p><ul><li><code>acc.declare</code> - This is a facility for easier determination of variables which are <code>acc declare</code>’d. This attribute is used on operations producing globals and on operations producing locals such as dialect specific <code>alloca</code>’s. Having this attribute is required in order to appear in a data mapping operation associated with any of the <code>acc.declare*</code> operations.</li><li><code>acc.declare_action</code> - Since the OpenACC specification allows declaration of variables that have yet to be allocated, this attribute is used at the allocation and deallocation points. More specifically, this attribute captures symbols of functions to be called to perform an action either pre-allocate, post-allocate, pre-deallocate, or post-deallocate. Calls to these functions should be materialized when lowering OpenACC semantics to ensure proper data actions are done after the allocation/deallocation.</li></ul><h2 id=openacc-transforms-and-analyses>OpenACC Transforms and Analyses <a class=headline-hash href=#openacc-transforms-and-analyses>¶</a></h2><p>The design goal for the <code>acc</code> dialect is to be friendly to MLIR optimization passes including CSE and LICM. Additionally, since it is designed to recover original clauses, it makes late verification and analysis possible in the MLIR framework outside of the frontend.</p><p>This section describes a few MLIR-level passes for which the <code>acc</code> dialect design should be friendly for. This section is currently solely outlining the possibilities intended by the design and not necessarily existing passes.</p><h3 id=verification>Verification <a class=headline-hash href=#verification>¶</a></h3><p>Since the OpenACC dialect is not lossy with regards to its representation, it is possible to do OpenACC language semantic checking at the MLIR-level. What follows is a list of various semantic checks needed.</p><p>This first list is required to be done in the frontend because the <code>acc</code> dialect operations must be valid when constructed:</p><ul><li>Ensure that only listed clauses are allowed for each directive.</li><li>Ensure that only listed modifiers are allowed for each clause.</li></ul><p>However, the following are semantic checks that can be done at the MLIR-level (either in a separate pass or as part of the operation verifier):</p><ul><li>Specify the validity checks that each modifier needs. (eg num_gangs may need a positive integer).</li><li>Ensure valid clause nesting.</li><li>Validate clause restrictions which cannot appear with others.</li><li>Validate that no conflicting clauses are used on variables.</li></ul><p>Note that some of these checks can be even more precise when done at the MLIR level because optimizations like inlining and constant propagation expose detail that wouldn’t have been visible in the frontend.</p><h3 id=implicit-data-attributes>Implicit Data Attributes <a class=headline-hash href=#implicit-data-attributes>¶</a></h3><p>The OpenACC specification includes a section on <code>2.6.2 Variables with Implicitly Determined Data Attributes</code>. What this section describes are the data actions that should be applied to a variable for which user did not specify a data action for. The action depends on the construct being used and also on the default clause. However, the point to note here is that variables which are live-in into the acc region must employ some data mapping so the data can be passed to accelerator.</p><p>One possible optimizations that affects data attributes needed is <code>Scalar Replacement of Aggregates (SROA)</code>. The <code>acc</code> dialect should not prevent this from happening on the source dialect.</p><p>Because it is intended to be possible to apply optimizations across an <code>acc</code> region, the analysis/transformation pass that applies the implicit data attributes should be run as late as possible - ideally right before any outlining process which uses the <code>acc</code> region body to create an accelerator procedure. It is expected that existing MLIR facilities, such as <code>mlir::Liveness</code> will work for the <code>acc</code> region and thus can be used to perform this analysis.</p><h3 id=redundant-clause-elimination>Redundant Clause Elimination <a class=headline-hash href=#redundant-clause-elimination>¶</a></h3><p>The data operations are modeled in a way where data entry operations look like loads and data exit operations look like stores. Thus these operations are intended to be optimized in the following ways:</p><ul><li>Be able to eliminate redundant operations such as when an <code>acc.copyin</code> dominates another.</li><li>Be able to hoist/sink such operations out of loops.</li></ul><h2 id=operations-toc>Operations TOC <a class=headline-hash href=#operations-toc>¶</a></h2><p><nav id=TableOfContents><ul><li><a href=#dialect-design-goals>Dialect Design Goals</a></li><li><a href=#operation-categories>Operation Categories</a><ul><li><a href=#data-operations>Data Operations</a></li><li><a href=#types>Types</a></li><li><a href=#recipes>Recipes</a></li><li><a href=#routine>Routine</a></li><li><a href=#declare>Declare</a></li></ul></li><li><a href=#openacc-transforms-and-analyses>OpenACC Transforms and Analyses</a><ul><li><a href=#verification>Verification</a></li><li><a href=#implicit-data-attributes>Implicit Data Attributes</a></li><li><a href=#redundant-clause-elimination>Redundant Clause Elimination</a></li></ul></li><li><a href=#operations-toc>Operations TOC</a></li><li><a href=#operations>Operations</a><ul><li><a href=#accatomiccapture-accatomiccaptureop><code>acc.atomic.capture</code> (acc::AtomicCaptureOp)</a></li><li><a href=#accatomicread-accatomicreadop><code>acc.atomic.read</code> (acc::AtomicReadOp)</a></li><li><a href=#accatomicupdate-accatomicupdateop><code>acc.atomic.update</code> (acc::AtomicUpdateOp)</a></li><li><a href=#accatomicwrite-accatomicwriteop><code>acc.atomic.write</code> (acc::AtomicWriteOp)</a></li><li><a href=#accattach-accattachop><code>acc.attach</code> (acc::AttachOp)</a></li><li><a href=#accbounds-accdataboundsop><code>acc.bounds</code> (acc::DataBoundsOp)</a></li><li><a href=#acccache-acccacheop><code>acc.cache</code> (acc::CacheOp)</a></li><li><a href=#acccopyin-acccopyinop><code>acc.copyin</code> (acc::CopyinOp)</a></li><li><a href=#acccopyout-acccopyoutop><code>acc.copyout</code> (acc::CopyoutOp)</a></li><li><a href=#acccreate-acccreateop><code>acc.create</code> (acc::CreateOp)</a></li><li><a href=#accdata-accdataop><code>acc.data</code> (acc::DataOp)</a></li><li><a href=#accdeclare-accdeclareop><code>acc.declare</code> (acc::DeclareOp)</a></li><li><a href=#accdeclare_device_resident-accdeclaredeviceresidentop><code>acc.declare_device_resident</code> (acc::DeclareDeviceResidentOp)</a></li><li><a href=#accdeclare_enter-accdeclareenterop><code>acc.declare_enter</code> (acc::DeclareEnterOp)</a></li><li><a href=#accdeclare_exit-accdeclareexitop><code>acc.declare_exit</code> (acc::DeclareExitOp)</a></li><li><a href=#accdeclare_link-accdeclarelinkop><code>acc.declare_link</code> (acc::DeclareLinkOp)</a></li><li><a href=#accdelete-accdeleteop><code>acc.delete</code> (acc::DeleteOp)</a></li><li><a href=#accdetach-accdetachop><code>acc.detach</code> (acc::DetachOp)</a></li><li><a href=#accdeviceptr-accdeviceptrop><code>acc.deviceptr</code> (acc::DevicePtrOp)</a></li><li><a href=#accenter_data-accenterdataop><code>acc.enter_data</code> (acc::EnterDataOp)</a></li><li><a href=#accexit_data-accexitdataop><code>acc.exit_data</code> (acc::ExitDataOp)</a></li><li><a href=#accfirstprivate-accfirstprivateop><code>acc.firstprivate</code> (acc::FirstprivateOp)</a></li><li><a href=#accfirstprivaterecipe-accfirstprivaterecipeop><code>acc.firstprivate.recipe</code> (acc::FirstprivateRecipeOp)</a></li><li><a href=#accgetdeviceptr-accgetdeviceptrop><code>acc.getdeviceptr</code> (acc::GetDevicePtrOp)</a></li><li><a href=#accglobal_ctor-accglobalconstructorop><code>acc.global_ctor</code> (acc::GlobalConstructorOp)</a></li><li><a href=#accglobal_dtor-accglobaldestructorop><code>acc.global_dtor</code> (acc::GlobalDestructorOp)</a></li><li><a href=#acchost_data-acchostdataop><code>acc.host_data</code> (acc::HostDataOp)</a></li><li><a href=#accinit-accinitop><code>acc.init</code> (acc::InitOp)</a></li><li><a href=#acckernels-acckernelsop><code>acc.kernels</code> (acc::KernelsOp)</a></li><li><a href=#accloop-accloopop><code>acc.loop</code> (acc::LoopOp)</a></li><li><a href=#accnocreate-accnocreateop><code>acc.nocreate</code> (acc::NoCreateOp)</a></li><li><a href=#accparallel-accparallelop><code>acc.parallel</code> (acc::ParallelOp)</a></li><li><a href=#accpresent-accpresentop><code>acc.present</code> (acc::PresentOp)</a></li><li><a href=#accprivate-accprivateop><code>acc.private</code> (acc::PrivateOp)</a></li><li><a href=#accprivaterecipe-accprivaterecipeop><code>acc.private.recipe</code> (acc::PrivateRecipeOp)</a></li><li><a href=#accreduction-accreductionop><code>acc.reduction</code> (acc::ReductionOp)</a></li><li><a href=#accreductionrecipe-accreductionrecipeop><code>acc.reduction.recipe</code> (acc::ReductionRecipeOp)</a></li><li><a href=#accroutine-accroutineop><code>acc.routine</code> (acc::RoutineOp)</a></li><li><a href=#accserial-accserialop><code>acc.serial</code> (acc::SerialOp)</a></li><li><a href=#accset-accsetop><code>acc.set</code> (acc::SetOp)</a></li><li><a href=#accshutdown-accshutdownop><code>acc.shutdown</code> (acc::ShutdownOp)</a></li><li><a href=#accterminator-accterminatorop><code>acc.terminator</code> (acc::TerminatorOp)</a></li><li><a href=#accupdate-accupdateop><code>acc.update</code> (acc::UpdateOp)</a></li><li><a href=#accupdate_device-accupdatedeviceop><code>acc.update_device</code> (acc::UpdateDeviceOp)</a></li><li><a href=#accupdate_host-accupdatehostop><code>acc.update_host</code> (acc::UpdateHostOp)</a></li><li><a href=#accuse_device-accusedeviceop><code>acc.use_device</code> (acc::UseDeviceOp)</a></li><li><a href=#accwait-accwaitop><code>acc.wait</code> (acc::WaitOp)</a></li><li><a href=#accyield-accyieldop><code>acc.yield</code> (acc::YieldOp)</a></li></ul></li><li><a href=#attributes-40>Attributes</a><ul><li><a href=#declareactionattr>DeclareActionAttr</a></li><li><a href=#declareattr>DeclareAttr</a></li><li><a href=#clausedefaultvalueattr>ClauseDefaultValueAttr</a></li><li><a href=#combinedconstructstypeattr>CombinedConstructsTypeAttr</a></li><li><a href=#constructattr>ConstructAttr</a></li><li><a href=#dataclauseattr>DataClauseAttr</a></li><li><a href=#devicetypeattr>DeviceTypeAttr</a></li><li><a href=#gangargtypeattr>GangArgTypeAttr</a></li><li><a href=#reductionoperatorattr>ReductionOperatorAttr</a></li><li><a href=#routineinfoattr>RoutineInfoAttr</a></li></ul></li><li><a href=#types-1>Types</a><ul><li><a href=#databoundstype>DataBoundsType</a></li><li><a href=#declaretokentype>DeclareTokenType</a></li></ul></li><li><a href=#enums>Enums</a><ul><li><a href=#clausedefaultvalue>ClauseDefaultValue</a></li><li><a href=#combinedconstructstype>CombinedConstructsType</a></li><li><a href=#construct>Construct</a></li><li><a href=#dataclause>DataClause</a></li><li><a href=#devicetype>DeviceType</a></li><li><a href=#gangargtype>GangArgType</a></li><li><a href=#reductionoperator>ReductionOperator</a></li></ul></li></ul></nav><h2 id=operations>Operations <a class=headline-hash href=#operations>¶</a></h2><p><a href=https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td>source</a></p><h3 id=accatomiccapture-accatomiccaptureop><code>acc.atomic.capture</code> (acc::AtomicCaptureOp) <a class=headline-hash href=#accatomiccapture-accatomiccaptureop>¶</a></h3><p><em>Performs an atomic capture</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.atomic.capture` $region attr-dict </code></pre><p>This operation performs an atomic capture.</p><p>The region has the following allowed forms:</p><pre tabindex=0><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 } </code></pre><p>Traits: <code>RecursiveMemoryEffects</code>, <code>SingleBlockImplicitTerminator<TerminatorOp></code>, <code>SingleBlock</code></p><p>Interfaces: <code>AtomicCaptureOpInterface</code></p><h3 id=accatomicread-accatomicreadop><code>acc.atomic.read</code> (acc::AtomicReadOp) <a class=headline-hash href=#accatomicread-accatomicreadop>¶</a></h3><p><em>Performs an atomic read</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.atomic.read` $v `=` $x `:` type($v) `,` type($x) `,` $element_type attr-dict </code></pre><p>This operation performs an atomic read.</p><p>The operand <code>x</code> is the address from where the value is atomically read. The operand <code>v</code> is the address where the value is stored after reading.</p><p>Interfaces: <code>AtomicReadOpInterface</code></p><h4 id=attributes>Attributes: <a class=headline-hash href=#attributes>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>element_type</code></td><td>::mlir::TypeAttr</td><td>any type attribute</td></tr></table><h4 id=operands>Operands: <a class=headline-hash href=#operands>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>x</code></td><td>pointer-like type</td></tr><tr><td style=text-align:center><code>v</code></td><td>pointer-like type</td></tr></tbody></table><h3 id=accatomicupdate-accatomicupdateop><code>acc.atomic.update</code> (acc::AtomicUpdateOp) <a class=headline-hash href=#accatomicupdate-accatomicupdateop>¶</a></h3><p><em>Performs an atomic update</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.atomic.update` $x `:` type($x) $region attr-dict </code></pre><p>This operation performs an atomic update.</p><p>The operand <code>x</code> is exactly the same as the operand <code>x</code> in the OpenACC Standard (OpenACC 3.3, section 2.12). It is the address of the variable that is being updated. <code>x</code> is atomically read/written.</p><p>The region describes how to update the value of <code>x</code>. It takes the value at <code>x</code> as an input and must yield the updated value. Only the update to <code>x</code> 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.</p><p>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 <code>atomicrmw</code> 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.</p><p>Traits: <code>RecursiveMemoryEffects</code>, <code>SingleBlockImplicitTerminator<YieldOp></code>, <code>SingleBlock</code></p><p>Interfaces: <code>AtomicUpdateOpInterface</code></p><h4 id=operands-1>Operands: <a class=headline-hash href=#operands-1>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>x</code></td><td>pointer-like type</td></tr></tbody></table><h3 id=accatomicwrite-accatomicwriteop><code>acc.atomic.write</code> (acc::AtomicWriteOp) <a class=headline-hash href=#accatomicwrite-accatomicwriteop>¶</a></h3><p><em>Performs an atomic write</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.atomic.write` $x `=` $expr `:` type($x) `,` type($expr) attr-dict </code></pre><p>This operation performs an atomic write.</p><p>The operand <code>x</code> is the address to where the <code>expr</code> is atomically written w.r.t. multiple threads. The evaluation of <code>expr</code> need not be atomic w.r.t. the write to address. In general, the type(x) must dereference to type(expr).</p><p>Interfaces: <code>AtomicWriteOpInterface</code></p><h4 id=operands-2>Operands: <a class=headline-hash href=#operands-2>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>x</code></td><td>pointer-like type</td></tr><tr><td style=text-align:center><code>expr</code></td><td>any type</td></tr></tbody></table><h3 id=accattach-accattachop><code>acc.attach</code> (acc::AttachOp) <a class=headline-hash href=#accattach-accattachop>¶</a></h3><p><em>Represents acc attach semantics which updates a pointer in device memory with the corresponding device address of the pointee.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.attach` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-1>Attributes: <a class=headline-hash href=#attributes-1>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-3>Operands: <a class=headline-hash href=#operands-3>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results>Results: <a class=headline-hash href=#results>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accbounds-accdataboundsop><code>acc.bounds</code> (acc::DataBoundsOp) <a class=headline-hash href=#accbounds-accdataboundsop>¶</a></h3><p><em>Represents normalized bounds information for acc data clause.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.bounds` oilist( `lowerbound` `(` $lowerbound `:` type($lowerbound) `)` | `upperbound` `(` $upperbound `:` type($upperbound) `)` | `extent` `(` $extent `:` type($extent) `)` | `stride` `(` $stride `:` type($stride) `)` | `startIdx` `(` $startIdx `:` type($startIdx) `)` ) attr-dict </code></pre><p>This operation is used to record bounds used in acc data clause in a normalized fashion (zero-based). This works well with the <code>PointerLikeType</code> requirement in data clauses - since a <code>lowerbound</code> of 0 means looking at data at the zero offset from pointer.</p><p>The operation must have an <code>upperbound</code> or <code>extent</code> (or both are allowed - but not checked for consistency). When the source language’s arrays are not zero-based, the <code>startIdx</code> must specify the zero-position index.</p><p>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.</p><p>C++:</p><pre tabindex=0><code>int array[10]; #pragma acc copy(array[1:9]) </code></pre><p>=></p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>bounds lb<span class=p>(</span><span class=m>1</span><span class=p>)</span> ub<span class=p>(</span><span class=m>9</span><span class=p>)</span> extent<span class=p>(</span><span class=m>9</span><span class=p>)</span> startIdx<span class=p>(</span><span class=m>0</span><span class=p>)</span> </span></span></code></pre></div><p>Fortran:</p><pre tabindex=0><code>integer :: array(1:10) !$acc copy(array(2:10)) </code></pre><p>=></p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>bounds lb<span class=p>(</span><span class=m>1</span><span class=p>)</span> ub<span class=p>(</span><span class=m>9</span><span class=p>)</span> extent<span class=p>(</span><span class=m>9</span><span class=p>)</span> startIdx<span class=p>(</span><span class=m>1</span><span class=p>)</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>NoMemoryEffect (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{}</code></p><h4 id=attributes-2>Attributes: <a class=headline-hash href=#attributes-2>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>strideInBytes</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr></table><h4 id=operands-4>Operands: <a class=headline-hash href=#operands-4>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>lowerbound</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>upperbound</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>extent</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>stride</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>startIdx</code></td><td>integer or index</td></tr></tbody></table><h4 id=results-1>Results: <a class=headline-hash href=#results-1>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>result</code></td><td>Type for representing acc data clause bounds information</td></tr></tbody></table><h3 id=acccache-acccacheop><code>acc.cache</code> (acc::CacheOp) <a class=headline-hash href=#acccache-acccacheop>¶</a></h3><p><em>Represents the cache directive that is associated with a loop.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.cache` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code>, <code>NoMemoryEffect (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{}</code></p><h4 id=attributes-3>Attributes: <a class=headline-hash href=#attributes-3>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-5>Operands: <a class=headline-hash href=#operands-5>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-2>Results: <a class=headline-hash href=#results-2>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=acccopyin-acccopyinop><code>acc.copyin</code> (acc::CopyinOp) <a class=headline-hash href=#acccopyin-acccopyinop>¶</a></h3><p><em>Represents copyin semantics for acc data clauses like acc copyin and acc copy.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.copyin` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-4>Attributes: <a class=headline-hash href=#attributes-4>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-6>Operands: <a class=headline-hash href=#operands-6>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-3>Results: <a class=headline-hash href=#results-3>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=acccopyout-acccopyoutop><code>acc.copyout</code> (acc::CopyoutOp) <a class=headline-hash href=#acccopyout-acccopyoutop>¶</a></h3><p><em>Represents acc copyout semantics - reverse of copyin.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.copyout` `accPtr` `(` $accPtr `:` type($accPtr) `)` (`bounds` `(` $bounds^ `)` )? (`async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType)^ `)`)? `to` `varPtr` `(` $varPtr `:` type($varPtr) `)` attr-dict </code></pre><ul><li><p><code>varPtr</code>: The address of variable to copy back to.</p><ul><li><code>accPtr</code>: The acc address of variable. This is the link from the data-entry operation used.</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p></li></ul><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-5>Attributes: <a class=headline-hash href=#attributes-5>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-7>Operands: <a class=headline-hash href=#operands-7>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h3 id=acccreate-acccreateop><code>acc.create</code> (acc::CreateOp) <a class=headline-hash href=#acccreate-acccreateop>¶</a></h3><p><em>Represents create semantics for acc data clauses like acc create and acc copyout.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.create` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-6>Attributes: <a class=headline-hash href=#attributes-6>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-8>Operands: <a class=headline-hash href=#operands-8>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-4>Results: <a class=headline-hash href=#results-4>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accdata-accdataop><code>acc.data</code> (acc::DataOp) <a class=headline-hash href=#accdata-accdataop>¶</a></h3><p><em>Data construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.data` oilist( `if` `(` $ifCond `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` | `wait` `` custom<WaitClause>($waitOperands, type($waitOperands), $waitOperandsDeviceType, $waitOperandsSegments, $hasWaitDevnum, $waitOnly) ) $region attr-dict-with-keyword </code></pre><p>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.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>data present<span class=p>(</span><span class=nv>%a</span><span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x10x</span><span class=k>f32</span><span class=p>>,</span> <span class=nv>%b</span><span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x10x</span><span class=k>f32</span><span class=p>>,</span> </span></span><span class=line><span class=cl> <span class=nv>%c</span><span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>,</span> <span class=nv>%d</span><span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// data region </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span></code></pre></div><p><code>async</code> and <code>wait</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code>, <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-7>Attributes: <a class=headline-hash href=#attributes-7>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>waitOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>waitOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>hasWaitDevnum</code></td><td>::mlir::ArrayAttr</td><td>1-bit boolean array attribute</td></tr><tr><td><code>waitOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>defaultAttr</code></td><td>::mlir::acc::ClauseDefaultValueAttr</td><td><details><summary>DefaultValue Clause</summary><p>Enum cases:</p><ul><li>present (<code>Present</code>)</li><li>none (<code>None</code>)</li></ul></details></td></tr></table><h4 id=operands-9>Operands: <a class=headline-hash href=#operands-9>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accdeclare-accdeclareop><code>acc.declare</code> (acc::DeclareOp) <a class=headline-hash href=#accdeclare-accdeclareop>¶</a></h3><p><em>Declare implicit region</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.declare` `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` $region attr-dict-with-keyword </code></pre><p>The “acc.declare” operation represents an implicit declare region in function (and subroutine in Fortran).</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl><span class=nv>%pa</span> <span class=p>=</span> acc<span class=p>.</span>present varPtr<span class=p>(</span><span class=nv>%a</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>-></span> <span class=kt>memref</span><span class=p><</span><span class=m>10x10x</span><span class=k>f32</span><span class=p>></span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>declare dataOperands<span class=p>(</span><span class=nv>%pa</span><span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// implicit region </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource}</code></p><h4 id=operands-10>Operands: <a class=headline-hash href=#operands-10>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accdeclare_device_resident-accdeclaredeviceresidentop><code>acc.declare_device_resident</code> (acc::DeclareDeviceResidentOp) <a class=headline-hash href=#accdeclare_device_resident-accdeclaredeviceresidentop>¶</a></h3><p><em>Represents acc declare device_resident semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.declare_device_resident` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-8>Attributes: <a class=headline-hash href=#attributes-8>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-11>Operands: <a class=headline-hash href=#operands-11>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-5>Results: <a class=headline-hash href=#results-5>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accdeclare_enter-accdeclareenterop><code>acc.declare_enter</code> (acc::DeclareEnterOp) <a class=headline-hash href=#accdeclare_enter-accdeclareenterop>¶</a></h3><p><em>Declare directive - entry to implicit data region</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.declare_enter` oilist( `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) attr-dict-with-keyword </code></pre><p>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”.</p><p>Example showing <code>acc declare create(a)</code>:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl><span class=nv>%0</span> <span class=p>=</span> acc<span class=p>.</span>create varPtr<span class=p>(</span><span class=nv>%a</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>-></span> <span class=p>!</span>llvm<span class=p>.</span>ptr </span></span><span class=line><span class=cl>acc<span class=p>.</span>declare_enter dataOperands<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> </span></span></code></pre></div><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=operands-12>Operands: <a class=headline-hash href=#operands-12>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h4 id=results-6>Results: <a class=headline-hash href=#results-6>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>token</code></td><td>declare token type</td></tr></tbody></table><h3 id=accdeclare_exit-accdeclareexitop><code>acc.declare_exit</code> (acc::DeclareExitOp) <a class=headline-hash href=#accdeclare_exit-accdeclareexitop>¶</a></h3><p><em>Declare directive - exit from implicit data region</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.declare_exit` oilist( `token` `(` $token `)` | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) attr-dict-with-keyword </code></pre><p>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”.</p><p>Example showing <code>acc declare device_resident(a)</code>:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl><span class=nv>%0</span> <span class=p>=</span> acc<span class=p>.</span>getdeviceptr varPtr<span class=p>(</span><span class=nv>%a</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>-></span> <span class=p>!</span>llvm<span class=p>.</span>ptr <span class=p>{</span><span class=nl>dataClause =</span> <span class=nv>#acc</span><span class=p><</span>data_clause declare_device_resident<span class=p>>}</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>declare_exit dataOperands<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>delete accPtr<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>{</span><span class=nl>dataClause =</span> <span class=nv>#acc</span><span class=p><</span>data_clause declare_device_resident<span class=p>>}</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=operands-13>Operands: <a class=headline-hash href=#operands-13>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>token</code></td><td>declare token type</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accdeclare_link-accdeclarelinkop><code>acc.declare_link</code> (acc::DeclareLinkOp) <a class=headline-hash href=#accdeclare_link-accdeclarelinkop>¶</a></h3><p><em>Represents acc declare link semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.declare_link` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-9>Attributes: <a class=headline-hash href=#attributes-9>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-14>Operands: <a class=headline-hash href=#operands-14>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-7>Results: <a class=headline-hash href=#results-7>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accdelete-accdeleteop><code>acc.delete</code> (acc::DeleteOp) <a class=headline-hash href=#accdelete-accdeleteop>¶</a></h3><p><em>Represents acc delete semantics - reverse of create.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.delete` `accPtr` `(` $accPtr `:` type($accPtr) `)` (`bounds` `(` $bounds^ `)` )? (`async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType)^ `)`)? attr-dict </code></pre><ul><li><code>accPtr</code>: The acc address of variable. This is the link from the data-entry operation used.</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-10>Attributes: <a class=headline-hash href=#attributes-10>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-15>Operands: <a class=headline-hash href=#operands-15>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h3 id=accdetach-accdetachop><code>acc.detach</code> (acc::DetachOp) <a class=headline-hash href=#accdetach-accdetachop>¶</a></h3><p><em>Represents acc detach semantics - reverse of attach.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.detach` `accPtr` `(` $accPtr `:` type($accPtr) `)` (`bounds` `(` $bounds^ `)` )? (`async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType)^ `)`)? attr-dict </code></pre><ul><li><code>accPtr</code>: The acc address of variable. This is the link from the data-entry operation used.</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-11>Attributes: <a class=headline-hash href=#attributes-11>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-16>Operands: <a class=headline-hash href=#operands-16>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h3 id=accdeviceptr-accdeviceptrop><code>acc.deviceptr</code> (acc::DevicePtrOp) <a class=headline-hash href=#accdeviceptr-accdeviceptrop>¶</a></h3><p><em>Specifies that the variable pointer is a device pointer.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.deviceptr` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-12>Attributes: <a class=headline-hash href=#attributes-12>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-17>Operands: <a class=headline-hash href=#operands-17>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-8>Results: <a class=headline-hash href=#results-8>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accenter_data-accenterdataop><code>acc.enter_data</code> (acc::EnterDataOp) <a class=headline-hash href=#accenter_data-accenterdataop>¶</a></h3><p><em>Enter data operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.enter_data` oilist( `if` `(` $ifCond `)` | `async` `(` $asyncOperand `:` type($asyncOperand) `)` | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)` | `wait` `(` $waitOperands `:` type($waitOperands) `)` | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) attr-dict-with-keyword </code></pre><p>The “acc.enter_data” operation represents the OpenACC enter data directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>enter_data create<span class=p>(</span><span class=nv>%d1</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> attributes <span class=p>{</span>async<span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-13>Attributes: <a class=headline-hash href=#attributes-13>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>async</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>wait</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-18>Operands: <a class=headline-hash href=#operands-18>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>asyncOperand</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>waitDevnum</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accexit_data-accexitdataop><code>acc.exit_data</code> (acc::ExitDataOp) <a class=headline-hash href=#accexit_data-accexitdataop>¶</a></h3><p><em>Exit data operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.exit_data` oilist( `if` `(` $ifCond `)` | `async` `(` $asyncOperand `:` type($asyncOperand) `)` | `wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)` | `wait` `(` $waitOperands `:` type($waitOperands) `)` | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) attr-dict-with-keyword </code></pre><p>The “acc.exit_data” operation represents the OpenACC exit data directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>exit_data delete<span class=p>(</span><span class=nv>%d1</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> attributes <span class=p>{</span>async<span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-14>Attributes: <a class=headline-hash href=#attributes-14>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>async</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>wait</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>finalize</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-19>Operands: <a class=headline-hash href=#operands-19>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>asyncOperand</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>waitDevnum</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accfirstprivate-accfirstprivateop><code>acc.firstprivate</code> (acc::FirstprivateOp) <a class=headline-hash href=#accfirstprivate-accfirstprivateop>¶</a></h3><p><em>Represents firstprivate semantic for the acc firstprivate clause.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.firstprivate` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-15>Attributes: <a class=headline-hash href=#attributes-15>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-20>Operands: <a class=headline-hash href=#operands-20>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-9>Results: <a class=headline-hash href=#results-9>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accfirstprivaterecipe-accfirstprivaterecipeop><code>acc.firstprivate.recipe</code> (acc::FirstprivateRecipeOp) <a class=headline-hash href=#accfirstprivaterecipe-accfirstprivaterecipeop>¶</a></h3><p><em>Privatization recipe</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.firstprivate.recipe` $sym_name `:` $type attr-dict-with-keyword `init` $initRegion `copy` $copyRegion (`destroy` $destroyRegion^)? </code></pre><p>Declares an OpenACC privatization recipe with copy of the initial value. The operation requires two mandatory regions and one optional.</p><ol><li>The initializer region specifies how to allocate and initialize a new private value. For example in Fortran, a derived-type might have a default initialization. The region has an argument that contains the value that need to be privatized. This is useful if the type is not known at compile time and the private value is needed to create its copy.</li><li>The copy region specifies how to copy the initial value to the newly created private value. It takes the initial value and the privatized value as arguments.</li><li>The destroy region specifies how to destruct the value when it reaches its end of life. It takes the privatized value as argument. It is optional.</li></ol><p>A single privatization recipe can be used for multiple operand if they have the same type and do not require a specific default initialization.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>firstprivate<span class=p>.</span>recipe <span class=nf>@privatization_f32</span> <span class=p>:</span> <span class=k>f32</span> init <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>f32</span><span class=p>):</span> </span></span><span class=line><span class=cl> <span class=c>// init region contains a sequence of operations to create and </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// initialize the copy if needed. It yields the create copy. </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> copy <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>f32</span><span class=p>,</span> <span class=nv>%1</span><span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>):</span> </span></span><span class=line><span class=cl> <span class=c>// copy region contains a sequence of operations to copy the initial value </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// of the firstprivate value to the newly created value. </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> destroy <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>f32</span><span class=p>)</span> </span></span><span class=line><span class=cl> <span class=c>// destroy region contains a sequences of operations to destruct the </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// created copy. </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span><span class=line><span class=cl> </span></span><span class=line><span class=cl><span class=c>// The privatization symbol is then used in the corresponding operation. </span></span></span><span class=line><span class=cl><span class=c></span>acc<span class=p>.</span>parallel firstprivate<span class=p>(</span><span class=nf>@privatization_f32</span> <span class=p>-></span> <span class=nv>%a</span> <span class=p>:</span> <span class=k>f32</span><span class=p>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>IsolatedFromAbove</code></p><p>Interfaces: <code>RecipeInterface</code>, <code>Symbol</code></p><h4 id=attributes-16>Attributes: <a class=headline-hash href=#attributes-16>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr><tr><td><code>type</code></td><td>::mlir::TypeAttr</td><td>any type attribute</td></tr></table><h3 id=accgetdeviceptr-accgetdeviceptrop><code>acc.getdeviceptr</code> (acc::GetDevicePtrOp) <a class=headline-hash href=#accgetdeviceptr-accgetdeviceptrop>¶</a></h3><p><em>Gets device address if variable exists on device.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.getdeviceptr` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>This operation is used to get the <code>accPtr</code> 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 <code>dataClause</code> argument that is any of the valid <code>mlir::acc::DataClause</code> entries. \</p><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-17>Attributes: <a class=headline-hash href=#attributes-17>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-21>Operands: <a class=headline-hash href=#operands-21>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-10>Results: <a class=headline-hash href=#results-10>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accglobal_ctor-accglobalconstructorop><code>acc.global_ctor</code> (acc::GlobalConstructorOp) <a class=headline-hash href=#accglobal_ctor-accglobalconstructorop>¶</a></h3><p><em>Used to hold construction operations associated with globals such as declare</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.global_ctor` $sym_name $region attr-dict-with-keyword </code></pre><p>The “acc.global_ctor” operation is used to capture OpenACC actions to apply on globals (such as <code>acc declare</code>) at the entry to the implicit data region. This operation is isolated and intended to be used in a module.</p><p>Example showing <code>declare create</code> of global:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>llvm<span class=p>.</span>mlir<span class=p>.</span>global external <span class=nf>@globalvar</span><span class=p>()</span> <span class=p>:</span> <span class=k>i32</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=nv>%0</span> <span class=p>=</span> llvm<span class=p>.</span>mlir<span class=p>.</span><span class=kt>constant</span><span class=p>(</span><span class=m>0</span> <span class=p>:</span> <span class=k>i32</span><span class=p>)</span> <span class=p>:</span> <span class=k>i32</span> </span></span><span class=line><span class=cl> llvm<span class=p>.</span><span class=kt>return</span> <span class=nv>%0</span> <span class=p>:</span> <span class=k>i32</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>global_ctor <span class=nf>@acc_constructor</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=nv>%0</span> <span class=p>=</span> llvm<span class=p>.</span>mlir<span class=p>.</span>addressof <span class=nf>@globalvar</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr </span></span><span class=line><span class=cl> <span class=nv>%1</span> <span class=p>=</span> acc<span class=p>.</span>create varPtr<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>-></span> <span class=p>!</span>llvm<span class=p>.</span>ptr </span></span><span class=line><span class=cl> acc<span class=p>.</span>declare_enter dataOperands<span class=p>(</span><span class=nv>%1</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>IsolatedFromAbove</code></p><p>Interfaces: <code>Symbol</code></p><h4 id=attributes-18>Attributes: <a class=headline-hash href=#attributes-18>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h3 id=accglobal_dtor-accglobaldestructorop><code>acc.global_dtor</code> (acc::GlobalDestructorOp) <a class=headline-hash href=#accglobal_dtor-accglobaldestructorop>¶</a></h3><p><em>Used to hold destruction operations associated with globals such as declare</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.global_dtor` $sym_name $region attr-dict-with-keyword </code></pre><p>The “acc.global_dtor” operation is used to capture OpenACC actions to apply on globals (such as <code>acc declare</code>) at the exit from the implicit data region. This operation is isolated and intended to be used in a module.</p><p>Example showing delete associated with <code>declare create</code> of global:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>llvm<span class=p>.</span>mlir<span class=p>.</span>global external <span class=nf>@globalvar</span><span class=p>()</span> <span class=p>:</span> <span class=k>i32</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=nv>%0</span> <span class=p>=</span> llvm<span class=p>.</span>mlir<span class=p>.</span><span class=kt>constant</span><span class=p>(</span><span class=m>0</span> <span class=p>:</span> <span class=k>i32</span><span class=p>)</span> <span class=p>:</span> <span class=k>i32</span> </span></span><span class=line><span class=cl> llvm<span class=p>.</span><span class=kt>return</span> <span class=nv>%0</span> <span class=p>:</span> <span class=k>i32</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>global_dtor <span class=nf>@acc_destructor</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=nv>%0</span> <span class=p>=</span> llvm<span class=p>.</span>mlir<span class=p>.</span>addressof <span class=nf>@globalvar</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr </span></span><span class=line><span class=cl> <span class=nv>%1</span> <span class=p>=</span> acc<span class=p>.</span>getdeviceptr varPtr<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>-></span> <span class=p>!</span>llvm<span class=p>.</span>ptr <span class=p>{</span><span class=nl>dataClause =</span> <span class=nv>#acc</span><span class=p><</span>data_clause create<span class=p>>}</span> </span></span><span class=line><span class=cl> acc<span class=p>.</span>declare_exit dataOperands<span class=p>(</span><span class=nv>%1</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> </span></span><span class=line><span class=cl> acc<span class=p>.</span>delete accPtr<span class=p>(</span><span class=nv>%1</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>{</span><span class=nl>dataClause =</span> <span class=nv>#acc</span><span class=p><</span>data_clause create<span class=p>>}</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>IsolatedFromAbove</code></p><p>Interfaces: <code>Symbol</code></p><h4 id=attributes-19>Attributes: <a class=headline-hash href=#attributes-19>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h3 id=acchost_data-acchostdataop><code>acc.host_data</code> (acc::HostDataOp) <a class=headline-hash href=#acchost_data-acchostdataop>¶</a></h3><p><em>Host_data construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.host_data` oilist( `if` `(` $ifCond `)` | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) $region attr-dict-with-keyword </code></pre><p>The “acc.host_data” operation represents the OpenACC host_data construct.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl><span class=nv>%0</span> <span class=p>=</span> acc<span class=p>.</span>use_device varPtr<span class=p>(</span><span class=nv>%a</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>-></span> <span class=p>!</span>llvm<span class=p>.</span>ptr </span></span><span class=line><span class=cl>acc<span class=p>.</span>host_data dataOperands<span class=p>(</span><span class=nv>%0</span> <span class=p>:</span> <span class=p>!</span>llvm<span class=p>.</span>ptr<span class=p>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-20>Attributes: <a class=headline-hash href=#attributes-20>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>ifPresent</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-22>Operands: <a class=headline-hash href=#operands-22>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accinit-accinitop><code>acc.init</code> (acc::InitOp) <a class=headline-hash href=#accinit-accinitop>¶</a></h3><p><em>Init operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.init` oilist(`device_num` `(` $deviceNumOperand `:` type($deviceNumOperand) `)` | `if` `(` $ifCond `)` ) attr-dict-with-keyword </code></pre><p>The “acc.init” operation represents the OpenACC init executable directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>init </span></span><span class=line><span class=cl>acc<span class=p>.</span>init device_num<span class=p>(</span><span class=nv>%dev1</span> <span class=p>:</span> <span class=k>i32</span><span class=p>)</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><h4 id=attributes-21>Attributes: <a class=headline-hash href=#attributes-21>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>device_types</code></td><td>::mlir::ArrayAttr</td><td>Device type attributes</td></tr></table><h4 id=operands-23>Operands: <a class=headline-hash href=#operands-23>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>deviceNumOperand</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr></tbody></table><h3 id=acckernels-acckernelsop><code>acc.kernels</code> (acc::KernelsOp) <a class=headline-hash href=#acckernels-acckernelsop>¶</a></h3><p><em>Kernels construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.kernels` ( `combined` `(` `loop` `)` $combined^)? oilist( `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` | `num_gangs` `(` custom<NumGangs>($numGangs, type($numGangs), $numGangsDeviceType, $numGangsSegments) `)` | `num_workers` `(` custom<DeviceTypeOperands>($numWorkers, type($numWorkers), $numWorkersDeviceType) `)` | `vector_length` `(` custom<DeviceTypeOperands>($vectorLength, type($vectorLength), $vectorLengthDeviceType) `)` | `wait` `` custom<WaitClause>($waitOperands, type($waitOperands), $waitOperandsDeviceType, $waitOperandsSegments, $hasWaitDevnum, $waitOnly) | `self` `(` $selfCond `)` | `if` `(` $ifCond `)` ) $region attr-dict-with-keyword </code></pre><p>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.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>kernels num_gangs<span class=p>(</span><span class=nv>%c10</span><span class=p>)</span> num_workers<span class=p>(</span><span class=nv>%c10</span><span class=p>)</span> </span></span><span class=line><span class=cl> private<span class=p>(</span><span class=nv>%c</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// kernels region </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span></code></pre></div><p><code>collapse</code>, <code>gang</code>, <code>worker</code>, <code>vector</code>, <code>seq</code>, <code>independent</code>, <code>auto</code> and <code>tile</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code>, <code>AutomaticAllocationScope</code>, <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>ComputeRegionOpInterface</code>, <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-22>Attributes: <a class=headline-hash href=#attributes-22>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>waitOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>waitOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>hasWaitDevnum</code></td><td>::mlir::ArrayAttr</td><td>1-bit boolean array attribute</td></tr><tr><td><code>waitOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>numGangsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>numGangsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>numWorkersDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>vectorLengthDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>selfAttr</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>defaultAttr</code></td><td>::mlir::acc::ClauseDefaultValueAttr</td><td><details><summary>DefaultValue Clause</summary><p>Enum cases:</p><ul><li>present (<code>Present</code>)</li><li>none (<code>None</code>)</li></ul></details></td></tr><tr><td><code>combined</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-24>Operands: <a class=headline-hash href=#operands-24>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>numGangs</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>numWorkers</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>vectorLength</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>selfCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accloop-accloopop><code>acc.loop</code> (acc::LoopOp) <a class=headline-hash href=#accloop-accloopop>¶</a></h3><p><em>Loop construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.loop` custom<CombinedConstructsLoop>($combined) oilist( `gang` `` custom<GangClause>($gangOperands, type($gangOperands), $gangOperandsArgType, $gangOperandsDeviceType, $gangOperandsSegments, $gang) | `worker` `` custom<DeviceTypeOperandsWithKeywordOnly>( $workerNumOperands, type($workerNumOperands), $workerNumOperandsDeviceType, $worker) | `vector` `` custom<DeviceTypeOperandsWithKeywordOnly>($vectorOperands, type($vectorOperands), $vectorOperandsDeviceType, $vector) | `private` `(` custom<SymOperandList>( $privateOperands, type($privateOperands), $privatizations) `)` | `tile` `(` custom<DeviceTypeOperandsWithSegment>($tileOperands, type($tileOperands), $tileOperandsDeviceType, $tileOperandsSegments) `)` | `reduction` `(` custom<SymOperandList>( $reductionOperands, type($reductionOperands), $reductionRecipes) `)` | `cache` `(` $cacheOperands `:` type($cacheOperands) `)` ) custom<LoopControl>($region, $lowerbound, type($lowerbound), $upperbound, type($upperbound), $step, type($step)) ( `(` type($results)^ `)` )? attr-dict-with-keyword </code></pre><p>The “acc.loop” operation represents the OpenACC loop construct. 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 <code>inclusive</code> attribute is set then the upper bound is included.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>loop gang<span class=p>()</span> <span class=kt>vector</span><span class=p>()</span> <span class=p>(</span><span class=nv>%arg3</span> <span class=p>:</span> <span class=k>index</span><span class=p>,</span> <span class=nv>%arg4</span> <span class=p>:</span> <span class=k>index</span><span class=p>,</span> <span class=nv>%arg5</span> <span class=p>:</span> <span class=k>index</span><span class=p>)</span> <span class=p>=</span> </span></span><span class=line><span class=cl> <span class=p>(</span><span class=nv>%c0</span><span class=p>,</span> <span class=nv>%c0</span><span class=p>,</span> <span class=nv>%c0</span> <span class=p>:</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>)</span> to </span></span><span class=line><span class=cl> <span class=p>(</span><span class=nv>%c10</span><span class=p>,</span> <span class=nv>%c10</span><span class=p>,</span> <span class=nv>%c10</span> <span class=p>:</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>)</span> step </span></span><span class=line><span class=cl> <span class=p>(</span><span class=nv>%c1</span><span class=p>,</span> <span class=nv>%c1</span><span class=p>,</span> <span class=nv>%c1</span> <span class=p>:</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>,</span> <span class=k>index</span><span class=p>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// Loop body </span></span></span><span class=line><span class=cl><span class=c></span> acc<span class=p>.</span>yield </span></span><span class=line><span class=cl><span class=p>}</span> attributes <span class=p>{</span> <span class=nl>collapse =</span> <span class=p>[</span><span class=m>3</span><span class=p>]</span> <span class=p>}</span> </span></span></code></pre></div><p><code>collapse</code>, <code>gang</code>, <code>worker</code>, <code>vector</code>, <code>seq</code>, <code>independent</code>, <code>auto</code> and <code>tile</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code>, <code>AutomaticAllocationScope</code>, <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>ComputeRegionOpInterface</code>, <code>LoopLikeOpInterface</code>, <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource}</code></p><h4 id=attributes-23>Attributes: <a class=headline-hash href=#attributes-23>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>inclusiveUpperbound</code></td><td>::mlir::DenseBoolArrayAttr</td><td>i1 dense array attribute</td></tr><tr><td><code>collapse</code></td><td>::mlir::ArrayAttr</td><td>64-bit integer array attribute</td></tr><tr><td><code>collapseDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>gangOperandsArgType</code></td><td>::mlir::ArrayAttr</td><td>gang arg type array attribute</td></tr><tr><td><code>gangOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>gangOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>workerNumOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>vectorOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>seq</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>independent</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>auto_</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>gang</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>worker</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>vector</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>tileOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>tileOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>privatizations</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>reductionRecipes</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>combined</code></td><td>::mlir::acc::CombinedConstructsTypeAttr</td><td><details><summary>Differentiate between combined constructs</summary><p>Enum cases:</p><ul><li>kernels_loop (<code>KernelsLoop</code>)</li><li>parallel_loop (<code>ParallelLoop</code>)</li><li>serial_loop (<code>SerialLoop</code>)</li></ul></details></td></tr></table><h4 id=operands-25>Operands: <a class=headline-hash href=#operands-25>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>lowerbound</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>upperbound</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>step</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>gangOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>workerNumOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>vectorOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>tileOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>cacheOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>privateOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>reductionOperands</code></td><td>variadic of any type</td></tr></tbody></table><h4 id=results-11>Results: <a class=headline-hash href=#results-11>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>results</code></td><td>variadic of any type</td></tr></tbody></table><h3 id=accnocreate-accnocreateop><code>acc.nocreate</code> (acc::NoCreateOp) <a class=headline-hash href=#accnocreate-accnocreateop>¶</a></h3><p><em>Represents acc no_create semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.nocreate` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-24>Attributes: <a class=headline-hash href=#attributes-24>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-26>Operands: <a class=headline-hash href=#operands-26>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-12>Results: <a class=headline-hash href=#results-12>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accparallel-accparallelop><code>acc.parallel</code> (acc::ParallelOp) <a class=headline-hash href=#accparallel-accparallelop>¶</a></h3><p><em>Parallel construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.parallel` ( `combined` `(` `loop` `)` $combined^)? oilist( `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` | `firstprivate` `(` custom<SymOperandList>($firstprivateOperands, type($firstprivateOperands), $firstprivatizations) `)` | `num_gangs` `(` custom<NumGangs>($numGangs, type($numGangs), $numGangsDeviceType, $numGangsSegments) `)` | `num_workers` `(` custom<DeviceTypeOperands>($numWorkers, type($numWorkers), $numWorkersDeviceType) `)` | `private` `(` custom<SymOperandList>( $privateOperands, type($privateOperands), $privatizations) `)` | `vector_length` `(` custom<DeviceTypeOperands>($vectorLength, type($vectorLength), $vectorLengthDeviceType) `)` | `wait` `` custom<WaitClause>($waitOperands, type($waitOperands), $waitOperandsDeviceType, $waitOperandsSegments, $hasWaitDevnum, $waitOnly) | `self` `(` $selfCond `)` | `if` `(` $ifCond `)` | `reduction` `(` custom<SymOperandList>( $reductionOperands, type($reductionOperands), $reductionRecipes) `)` ) $region attr-dict-with-keyword </code></pre><p>The “acc.parallel” operation represents a parallel construct block. It has one region to be executed in parallel on the current device.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>parallel num_gangs<span class=p>(</span><span class=nv>%c10</span><span class=p>)</span> num_workers<span class=p>(</span><span class=nv>%c10</span><span class=p>)</span> </span></span><span class=line><span class=cl> private<span class=p>(</span><span class=nv>%c</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// parallel region </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span></code></pre></div><p><code>async</code>, <code>wait</code>, <code>num_gangs</code>, <code>num_workers</code> and <code>vector_length</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code>, <code>AutomaticAllocationScope</code>, <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>ComputeRegionOpInterface</code>, <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-25>Attributes: <a class=headline-hash href=#attributes-25>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>waitOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>waitOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>hasWaitDevnum</code></td><td>::mlir::ArrayAttr</td><td>1-bit boolean array attribute</td></tr><tr><td><code>waitOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>numGangsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>numGangsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>numWorkersDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>vectorLengthDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>selfAttr</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>reductionRecipes</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>privatizations</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>firstprivatizations</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>defaultAttr</code></td><td>::mlir::acc::ClauseDefaultValueAttr</td><td><details><summary>DefaultValue Clause</summary><p>Enum cases:</p><ul><li>present (<code>Present</code>)</li><li>none (<code>None</code>)</li></ul></details></td></tr><tr><td><code>combined</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-27>Operands: <a class=headline-hash href=#operands-27>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>numGangs</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>numWorkers</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>vectorLength</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>selfCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>reductionOperands</code></td><td>variadic of any type</td></tr><tr><td style=text-align:center><code>privateOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>firstprivateOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accpresent-accpresentop><code>acc.present</code> (acc::PresentOp) <a class=headline-hash href=#accpresent-accpresentop>¶</a></h3><p><em>Specifies that the variable is already present on device.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.present` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-26>Attributes: <a class=headline-hash href=#attributes-26>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-28>Operands: <a class=headline-hash href=#operands-28>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-13>Results: <a class=headline-hash href=#results-13>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accprivate-accprivateop><code>acc.private</code> (acc::PrivateOp) <a class=headline-hash href=#accprivate-accprivateop>¶</a></h3><p><em>Represents private semantics for acc private clause.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.private` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-27>Attributes: <a class=headline-hash href=#attributes-27>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-29>Operands: <a class=headline-hash href=#operands-29>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-14>Results: <a class=headline-hash href=#results-14>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accprivaterecipe-accprivaterecipeop><code>acc.private.recipe</code> (acc::PrivateRecipeOp) <a class=headline-hash href=#accprivaterecipe-accprivaterecipeop>¶</a></h3><p><em>Privatization recipe</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.private.recipe` $sym_name `:` $type attr-dict-with-keyword `init` $initRegion (`destroy` $destroyRegion^)? </code></pre><p>Declares an OpenACC privatization recipe. The operation requires one mandatory and one optional region.</p><ol><li>The initializer region specifies how to allocate and initialize a new private value. For example in Fortran, a derived-type might have a default initialization. The region has an argument that contains the value that need to be privatized. This is useful if the type is not known at compile time and the private value is needed to create its copy.</li><li>The destroy region specifies how to destruct the value when it reaches its end of life. It takes the privatized value as argument.</li></ol><p>A single privatization recipe can be used for multiple operand if they have the same type and do not require a specific default initialization.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>private<span class=p>.</span>recipe <span class=nf>@privatization_f32</span> <span class=p>:</span> <span class=k>f32</span> init <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>f32</span><span class=p>):</span> </span></span><span class=line><span class=cl> <span class=c>// init region contains a sequence of operations to create and </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// initialize the copy if needed. It yields the create copy. </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> destroy <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>f32</span><span class=p>)</span> </span></span><span class=line><span class=cl> <span class=c>// destroy region contains a sequences of operations to destruct the </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// created copy. </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span><span class=line><span class=cl> </span></span><span class=line><span class=cl><span class=c>// The privatization symbol is then used in the corresponding operation. </span></span></span><span class=line><span class=cl><span class=c></span>acc<span class=p>.</span>parallel private<span class=p>(</span><span class=nf>@privatization_f32</span> <span class=p>-></span> <span class=nv>%a</span> <span class=p>:</span> <span class=k>f32</span><span class=p>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>Traits: <code>IsolatedFromAbove</code></p><p>Interfaces: <code>RecipeInterface</code>, <code>Symbol</code></p><h4 id=attributes-28>Attributes: <a class=headline-hash href=#attributes-28>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr><tr><td><code>type</code></td><td>::mlir::TypeAttr</td><td>any type attribute</td></tr></table><h3 id=accreduction-accreductionop><code>acc.reduction</code> (acc::ReductionOp) <a class=headline-hash href=#accreduction-accreductionop>¶</a></h3><p><em>Represents reduction semantics for acc reduction clause.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.reduction` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-29>Attributes: <a class=headline-hash href=#attributes-29>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-30>Operands: <a class=headline-hash href=#operands-30>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-15>Results: <a class=headline-hash href=#results-15>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accreductionrecipe-accreductionrecipeop><code>acc.reduction.recipe</code> (acc::ReductionRecipeOp) <a class=headline-hash href=#accreductionrecipe-accreductionrecipeop>¶</a></h3><p><em>Reduction recipe</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.reduction.recipe` $sym_name `:` $type attr-dict-with-keyword `reduction_operator` $reductionOperator `init` $initRegion `combiner` $combinerRegion </code></pre><p>Declares an OpenACC reduction recipe. The operation requires two mandatory regions.</p><ol><li>The initializer region specifies how to initialize the local reduction value. The region has a first argument that contains the value of the reduction accumulator at the start of the reduction. It is expected to <code>acc.yield</code> the new value. Extra arguments can be added to deal with dynamic arrays.</li><li>The reduction region contains a sequences of operations to combine two values of the reduction type into one. It has at least two arguments and it is expected to <code>acc.yield</code> the combined value. Extra arguments can be added to deal with dynamic arrays.</li></ol><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>reduction<span class=p>.</span>recipe <span class=nf>@reduction_add_i64</span> <span class=p>:</span> <span class=k>i64</span> reduction_operator<span class=p><</span>add<span class=p>></span> init <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>i64</span><span class=p>):</span> </span></span><span class=line><span class=cl> <span class=c>// init region contains a sequence of operations to initialize the local </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// reduction value as specified in 2.5.15 </span></span></span><span class=line><span class=cl><span class=c></span> <span class=nv>%c0</span> <span class=p>=</span> arith<span class=p>.</span><span class=kt>constant</span> <span class=m>0</span> <span class=p>:</span> <span class=k>i64</span> </span></span><span class=line><span class=cl> acc<span class=p>.</span>yield <span class=nv>%c0</span> <span class=p>:</span> <span class=k>i64</span> </span></span><span class=line><span class=cl><span class=p>}</span> combiner <span class=p>{</span> </span></span><span class=line><span class=cl><span class=nl>^bb0</span><span class=p>(</span><span class=nv>%0</span><span class=p>:</span> <span class=k>i64</span><span class=p>,</span> <span class=nv>%1</span><span class=p>:</span> <span class=k>i64</span><span class=p>)</span> </span></span><span class=line><span class=cl> <span class=c>// combiner region contains a sequence of operations to combine </span></span></span><span class=line><span class=cl><span class=c></span> <span class=c>// two values into one. </span></span></span><span class=line><span class=cl><span class=c></span> <span class=nv>%2</span> <span class=p>=</span> arith<span class=p>.</span>addi <span class=nv>%0</span><span class=p>,</span> <span class=nv>%1</span> <span class=p>:</span> <span class=k>i64</span> </span></span><span class=line><span class=cl> acc<span class=p>.</span>yield <span class=nv>%2</span> <span class=p>:</span> <span class=k>i64</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span><span class=line><span class=cl> </span></span><span class=line><span class=cl><span class=c>// The reduction symbol is then used in the corresponding operation. </span></span></span><span class=line><span class=cl><span class=c></span>acc<span class=p>.</span>parallel reduction<span class=p>(</span><span class=nf>@reduction_add_i64</span> <span class=p>-></span> <span class=nv>%a</span> <span class=p>:</span> <span class=k>i64</span><span class=p>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span></code></pre></div><p>The following table lists the valid operators and the initialization values according to OpenACC 3.3:</p><p>|————————————————|</p><table><thead><tr><th>C/C++</th><th>Fortran</th></tr></thead><tbody><tr><td>operator</td><td>init value</td></tr><tr><td>+</td><td>0</td></tr><tr><td>*</td><td>1</td></tr><tr><td>max</td><td>least</td></tr><tr><td>min</td><td>largest</td></tr><tr><td>&</td><td>~0</td></tr><tr><td></td><td></td></tr><tr><td>^</td><td>0</td></tr><tr><td>&&</td><td>1</td></tr><tr><td></td><td></td></tr><tr><td></td><td></td></tr><tr><td></td><td></td></tr><tr><td>————————————————-</td><td></td></tr></tbody></table><p>Traits: <code>IsolatedFromAbove</code></p><p>Interfaces: <code>RecipeInterface</code>, <code>Symbol</code></p><h4 id=attributes-30>Attributes: <a class=headline-hash href=#attributes-30>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr><tr><td><code>type</code></td><td>::mlir::TypeAttr</td><td>any type attribute</td></tr><tr><td><code>reductionOperator</code></td><td>::mlir::acc::ReductionOperatorAttr</td><td><details><summary>built-in reduction operations supported by OpenACC</summary><p>Enum cases:</p><ul><li>add (<code>AccAdd</code>)</li><li>mul (<code>AccMul</code>)</li><li>max (<code>AccMax</code>)</li><li>min (<code>AccMin</code>)</li><li>iand (<code>AccIand</code>)</li><li>ior (<code>AccIor</code>)</li><li>xor (<code>AccXor</code>)</li><li>eqv (<code>AccEqv</code>)</li><li>neqv (<code>AccNeqv</code>)</li><li>land (<code>AccLand</code>)</li><li>lor (<code>AccLor</code>)</li></ul></details></td></tr></table><h3 id=accroutine-accroutineop><code>acc.routine</code> (acc::RoutineOp) <a class=headline-hash href=#accroutine-accroutineop>¶</a></h3><p><em>Acc routine operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.routine` $sym_name `func` `(` $func_name `)` oilist ( `bind` `(` custom<BindName>($bindName, $bindNameDeviceType) `)` | `gang` `` custom<RoutineGangClause>($gang, $gangDim, $gangDimDeviceType) | `worker` custom<DeviceTypeArrayAttr>($worker) | `vector` custom<DeviceTypeArrayAttr>($vector) | `seq` custom<DeviceTypeArrayAttr>($seq) | `nohost` $nohost | `implicit` $implicit ) attr-dict-with-keyword </code></pre><p>The <code>acc.routine</code> 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 <code>RoutineInfoAttr</code>.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl><span class=kt>func</span><span class=p>.</span><span class=kt>func</span> <span class=nf>@acc_func</span><span class=p>(</span><span class=nv>%a</span> <span class=p>:</span> <span class=k>i64</span><span class=p>)</span> <span class=p>-></span> <span class=p>()</span> attributes </span></span><span class=line><span class=cl> <span class=p>{</span><span class=nl>acc.routine_info =</span> <span class=nv>#acc.routine_info</span><span class=p><[</span><span class=nf>@acc_func_rout1</span><span class=p>]>}</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=kt>return</span> </span></span><span class=line><span class=cl><span class=p>}</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>routine <span class=nf>@acc_func_rout1</span> <span class=kt>func</span><span class=p>(</span><span class=nf>@acc_func</span><span class=p>)</span> gang </span></span></code></pre></div><p><code>bind</code>, <code>gang</code>, <code>worker</code>, <code>vector</code> and <code>seq</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>IsolatedFromAbove</code></p><h4 id=attributes-31>Attributes: <a class=headline-hash href=#attributes-31>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr><tr><td><code>func_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr><tr><td><code>bindName</code></td><td>::mlir::ArrayAttr</td><td>string array attribute</td></tr><tr><td><code>bindNameDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>worker</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>vector</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>seq</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>nohost</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>gang</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>gangDim</code></td><td>::mlir::ArrayAttr</td><td>64-bit integer array attribute</td></tr><tr><td><code>gangDimDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr></table><h3 id=accserial-accserialop><code>acc.serial</code> (acc::SerialOp) <a class=headline-hash href=#accserial-accserialop>¶</a></h3><p><em>Serial construct</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.serial` ( `combined` `(` `loop` `)` $combined^)? oilist( `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` | `firstprivate` `(` custom<SymOperandList>($firstprivateOperands, type($firstprivateOperands), $firstprivatizations) `)` | `private` `(` custom<SymOperandList>( $privateOperands, type($privateOperands), $privatizations) `)` | `wait` `` custom<WaitClause>($waitOperands, type($waitOperands), $waitOperandsDeviceType, $waitOperandsSegments, $hasWaitDevnum, $waitOnly) | `self` `(` $selfCond `)` | `if` `(` $ifCond `)` | `reduction` `(` custom<SymOperandList>( $reductionOperands, type($reductionOperands), $reductionRecipes) `)` ) $region attr-dict-with-keyword </code></pre><p>The “acc.serial” operation represents a serial construct block. It has one region to be executed in serial on the current device.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>serial private<span class=p>(</span><span class=nv>%c</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> <span class=p>{</span> </span></span><span class=line><span class=cl> <span class=c>// serial region </span></span></span><span class=line><span class=cl><span class=c></span><span class=p>}</span> </span></span></code></pre></div><p><code>async</code> and <code>wait</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code>, <code>AutomaticAllocationScope</code>, <code>RecursiveMemoryEffects</code></p><p>Interfaces: <code>ComputeRegionOpInterface</code>, <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-32>Attributes: <a class=headline-hash href=#attributes-32>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>waitOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>waitOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>hasWaitDevnum</code></td><td>::mlir::ArrayAttr</td><td>1-bit boolean array attribute</td></tr><tr><td><code>waitOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>selfAttr</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr><tr><td><code>reductionRecipes</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>privatizations</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>firstprivatizations</code></td><td>::mlir::ArrayAttr</td><td>symbol ref array attribute</td></tr><tr><td><code>defaultAttr</code></td><td>::mlir::acc::ClauseDefaultValueAttr</td><td><details><summary>DefaultValue Clause</summary><p>Enum cases:</p><ul><li>present (<code>Present</code>)</li><li>none (<code>None</code>)</li></ul></details></td></tr><tr><td><code>combined</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-31>Operands: <a class=headline-hash href=#operands-31>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>selfCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>reductionOperands</code></td><td>variadic of any type</td></tr><tr><td style=text-align:center><code>privateOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>firstprivateOperands</code></td><td>variadic of PointerLikeType instance</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accset-accsetop><code>acc.set</code> (acc::SetOp) <a class=headline-hash href=#accset-accsetop>¶</a></h3><p><em>Set operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.set` oilist(`default_async` `(` $defaultAsync `:` type($defaultAsync) `)` | `device_num` `(` $deviceNum `:` type($deviceNum) `)` | `if` `(` $ifCond `)` ) attr-dict-with-keyword </code></pre><p>The “acc.set” operation represents the OpenACC set directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>set device_num<span class=p>(</span><span class=nv>%dev1</span> <span class=p>:</span> <span class=k>i32</span><span class=p>)</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-33>Attributes: <a class=headline-hash href=#attributes-33>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>device_type</code></td><td>::mlir::acc::DeviceTypeAttr</td><td><details><summary>built-in device type supported by OpenACC</summary><p>Enum cases:</p><ul><li>none (<code>None</code>)</li><li>star (<code>Star</code>)</li><li>default (<code>Default</code>)</li><li>host (<code>Host</code>)</li><li>multicore (<code>Multicore</code>)</li><li>nvidia (<code>Nvidia</code>)</li><li>radeon (<code>Radeon</code>)</li></ul></details></td></tr></table><h4 id=operands-32>Operands: <a class=headline-hash href=#operands-32>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>defaultAsync</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>deviceNum</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr></tbody></table><h3 id=accshutdown-accshutdownop><code>acc.shutdown</code> (acc::ShutdownOp) <a class=headline-hash href=#accshutdown-accshutdownop>¶</a></h3><p><em>Shutdown operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.shutdown` oilist(`device_num` `(` $deviceNumOperand `:` type($deviceNumOperand) `)` |`if` `(` $ifCond `)` ) attr-dict-with-keyword </code></pre><p>The “acc.shutdown” operation represents the OpenACC shutdown executable directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>shutdown </span></span><span class=line><span class=cl>acc<span class=p>.</span>shutdown device_num<span class=p>(</span><span class=nv>%dev1</span> <span class=p>:</span> <span class=k>i32</span><span class=p>)</span> </span></span></code></pre></div><p>Traits: <code>AttrSizedOperandSegments</code></p><h4 id=attributes-34>Attributes: <a class=headline-hash href=#attributes-34>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>device_types</code></td><td>::mlir::ArrayAttr</td><td>Device type attributes</td></tr></table><h4 id=operands-33>Operands: <a class=headline-hash href=#operands-33>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>deviceNumOperand</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr></tbody></table><h3 id=accterminator-accterminatorop><code>acc.terminator</code> (acc::TerminatorOp) <a class=headline-hash href=#accterminator-accterminatorop>¶</a></h3><p><em>Generic terminator for OpenACC regions</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.terminator` attr-dict </code></pre><p>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.</p><p>Traits: <code>AlwaysSpeculatableImplTrait</code>, <code>Terminator</code></p><p>Interfaces: <code>ConditionallySpeculatable</code>, <code>NoMemoryEffect (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{}</code></p><h3 id=accupdate-accupdateop><code>acc.update</code> (acc::UpdateOp) <a class=headline-hash href=#accupdate-accupdateop>¶</a></h3><p><em>Update operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.update` oilist( `if` `(` $ifCond `)` | `async` `` custom<DeviceTypeOperandsWithKeywordOnly>( $asyncOperands, type($asyncOperands), $asyncOperandsDeviceType, $async) | `wait` `` custom<WaitClause>($waitOperands, type($waitOperands), $waitOperandsDeviceType, $waitOperandsSegments, $hasWaitDevnum, $waitOnly) | `dataOperands` `(` $dataClauseOperands `:` type($dataClauseOperands) `)` ) attr-dict-with-keyword </code></pre><p>The <code>acc.update</code> operation represents the OpenACC update executable directive. As host and self clauses are synonyms, any operands for host and self are add to $hostOperands.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>update device<span class=p>(</span><span class=nv>%d1</span> <span class=p>:</span> <span class=kt>memref</span><span class=p><</span><span class=m>10x</span><span class=k>f32</span><span class=p>>)</span> attributes <span class=p>{</span>async<span class=p>}</span> </span></span></code></pre></div><p><code>async</code> and <code>wait</code> operands are supported with <code>device_type</code> information. They should only be accessed by the extra provided getters. If modified, the corresponding <code>device_type</code> attributes must be modified as well.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Write on ::mlir::acc::ConstructResource, MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-35>Attributes: <a class=headline-hash href=#attributes-35>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>async</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>waitOperandsSegments</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr><tr><td><code>waitOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>hasWaitDevnum</code></td><td>::mlir::ArrayAttr</td><td>1-bit boolean array attribute</td></tr><tr><td><code>waitOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>ifPresent</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-34>Operands: <a class=headline-hash href=#operands-34>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>dataClauseOperands</code></td><td>variadic of PointerLikeType instance</td></tr></tbody></table><h3 id=accupdate_device-accupdatedeviceop><code>acc.update_device</code> (acc::UpdateDeviceOp) <a class=headline-hash href=#accupdate_device-accupdatedeviceop>¶</a></h3><p><em>Represents acc update device semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.update_device` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code></p><h4 id=attributes-36>Attributes: <a class=headline-hash href=#attributes-36>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-35>Operands: <a class=headline-hash href=#operands-35>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-16>Results: <a class=headline-hash href=#results-16>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accupdate_host-accupdatehostop><code>acc.update_host</code> (acc::UpdateHostOp) <a class=headline-hash href=#accupdate_host-accupdatehostop>¶</a></h3><p><em>Represents acc update host semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.update_host` `accPtr` `(` $accPtr `:` type($accPtr) `)` (`bounds` `(` $bounds^ `)` )? (`async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType)^ `)`)? `to` `varPtr` `(` $varPtr `:` type($varPtr) `)` attr-dict </code></pre><ul><li><p><code>varPtr</code>: The address of variable to copy back to.</p><ul><li><code>accPtr</code>: The acc address of variable. This is the link from the data-entry operation used.</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p></li></ul><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters, MemoryEffects::Write on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-37>Attributes: <a class=headline-hash href=#attributes-37>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-36>Operands: <a class=headline-hash href=#operands-36>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h3 id=accuse_device-accusedeviceop><code>acc.use_device</code> (acc::UseDeviceOp) <a class=headline-hash href=#accuse_device-accusedeviceop>¶</a></h3><p><em>Represents acc use_device semantics.</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.use_device` `varPtr` `(` $varPtr `:` type($varPtr) `)` oilist( `varPtrPtr` `(` $varPtrPtr `:` type($varPtrPtr) `)` | `bounds` `(` $bounds `)` | `async` `(` custom<DeviceTypeOperands>($asyncOperands, type($asyncOperands), $asyncOperandsDeviceType) `)` ) `->` type($accPtr) attr-dict </code></pre><p>Description of arguments:</p><ul><li><code>varPtr</code>: The address of variable to copy.</li><li><code>varPtrPtr</code>: Specifies the address of varPtr - only used when the variable copied is a field in a struct. This is important for OpenACC due to implicit attach semantics on data clauses (2.6.4).</li><li><code>bounds</code>: 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.</li><li><code>asyncOperands</code> and <code>asyncOperandsDeviceType</code>: pair-wise lists of the async clause values associated with device_type’s.</li><li><code>asyncOnly</code>: 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).</li><li><code>dataClause</code>: 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 <code>acc.copyin</code> and <code>acc.copyout</code> operations, but both have dataClause that specifies <code>acc_copy</code> in this field.</li><li><code>structured</code>: 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).</li><li><code>implicit</code>: Whether this is an implicitly generated operation, such as copies done to satisfy “Variables with Implicitly Determined Data Attributes” in 2.6.2.</li><li><code>name</code>: Holds the name of variable as specified in user clause (including bounds).</li></ul><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><p>Interfaces: <code>MemoryEffectOpInterface (MemoryEffectOpInterface)</code></p><p>Effects: <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::CurrentDeviceIdResource}</code>, <code>MemoryEffects::Effect{MemoryEffects::Read on ::mlir::acc::RuntimeCounters}</code></p><h4 id=attributes-38>Attributes: <a class=headline-hash href=#attributes-38>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>asyncOperandsDeviceType</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>asyncOnly</code></td><td>::mlir::ArrayAttr</td><td>device type array attribute</td></tr><tr><td><code>dataClause</code></td><td>::mlir::acc::DataClauseAttr</td><td><details><summary>data clauses supported by OpenACC</summary><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul></details></td></tr><tr><td><code>structured</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>implicit</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr><tr><td><code>name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr></table><h4 id=operands-37>Operands: <a class=headline-hash href=#operands-37>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>varPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>varPtrPtr</code></td><td>PointerLikeType instance</td></tr><tr><td style=text-align:center><code>bounds</code></td><td>variadic of Type for representing acc data clause bounds information</td></tr><tr><td style=text-align:center><code>asyncOperands</code></td><td>variadic of integer or index</td></tr></tbody></table><h4 id=results-17>Results: <a class=headline-hash href=#results-17>¶</a></h4><table><thead><tr><th style=text-align:center>Result</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>accPtr</code></td><td>PointerLikeType instance</td></tr></tbody></table><h3 id=accwait-accwaitop><code>acc.wait</code> (acc::WaitOp) <a class=headline-hash href=#accwait-accwaitop>¶</a></h3><p><em>Wait operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.wait` ( `(` $waitOperands^ `:` type($waitOperands) `)` )? oilist(`async` `(` $asyncOperand `:` type($asyncOperand) `)` |`wait_devnum` `(` $waitDevnum `:` type($waitDevnum) `)` |`if` `(` $ifCond `)` ) attr-dict-with-keyword </code></pre><p>The “acc.wait” operation represents the OpenACC wait executable directive.</p><p>Example:</p><div class=highlight><pre tabindex=0 class=chroma><code class=language-mlir data-lang=mlir><span class=line><span class=cl>acc<span class=p>.</span>wait<span class=p>(</span><span class=nv>%value1</span><span class=p>:</span> <span class=k>index</span><span class=p>)</span> </span></span><span class=line><span class=cl>acc<span class=p>.</span>wait<span class=p>()</span> async<span class=p>(</span><span class=nv>%async1</span><span class=p>:</span> <span class=k>i32</span><span class=p>)</span> </span></span></code></pre></div><p>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.</p><p>Traits: <code>AttrSizedOperandSegments</code></p><h4 id=attributes-39>Attributes: <a class=headline-hash href=#attributes-39>¶</a></h4><table><tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr><tr><td><code>async</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr></table><h4 id=operands-38>Operands: <a class=headline-hash href=#operands-38>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>waitOperands</code></td><td>variadic of integer or index</td></tr><tr><td style=text-align:center><code>asyncOperand</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>waitDevnum</code></td><td>integer or index</td></tr><tr><td style=text-align:center><code>ifCond</code></td><td>1-bit signless integer</td></tr></tbody></table><h3 id=accyield-accyieldop><code>acc.yield</code> (acc::YieldOp) <a class=headline-hash href=#accyield-accyieldop>¶</a></h3><p><em>Acc yield and termination operation</em></p><p>Syntax:</p><pre tabindex=0><code>operation ::= `acc.yield` attr-dict ($operands^ `:` type($operands))? </code></pre><p><code>acc.yield</code> 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.</p><p>Traits: <code>AlwaysSpeculatableImplTrait</code>, <code>HasParent<FirstprivateRecipeOp, LoopOp, ParallelOp, PrivateRecipeOp,ReductionRecipeOp, SerialOp, AtomicUpdateOp></code>, <code>ReturnLike</code>, <code>Terminator</code></p><p>Interfaces: <code>ConditionallySpeculatable</code>, <code>NoMemoryEffect (MemoryEffectOpInterface)</code>, <code>RegionBranchTerminatorOpInterface</code></p><p>Effects: <code>MemoryEffects::Effect{}</code></p><h4 id=operands-39>Operands: <a class=headline-hash href=#operands-39>¶</a></h4><table><thead><tr><th style=text-align:center>Operand</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center><code>operands</code></td><td>variadic of any type</td></tr></tbody></table><h2 id=attributes-40>Attributes <a class=headline-hash href=#attributes-40>¶</a></h2><h3 id=declareactionattr>DeclareActionAttr <a class=headline-hash href=#declareactionattr>¶</a></h3><p>Syntax:</p><pre tabindex=0><code>#acc.declare_action< SymbolRefAttr, # preAlloc SymbolRefAttr, # postAlloc SymbolRefAttr, # preDealloc SymbolRefAttr # postDealloc > </code></pre><h4 id=parameters>Parameters: <a class=headline-hash href=#parameters>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>preAlloc</td><td style=text-align:center><code>SymbolRefAttr</code></td><td></td></tr><tr><td style=text-align:center>postAlloc</td><td style=text-align:center><code>SymbolRefAttr</code></td><td></td></tr><tr><td style=text-align:center>preDealloc</td><td style=text-align:center><code>SymbolRefAttr</code></td><td></td></tr><tr><td style=text-align:center>postDealloc</td><td style=text-align:center><code>SymbolRefAttr</code></td><td></td></tr></tbody></table><h3 id=declareattr>DeclareAttr <a class=headline-hash href=#declareattr>¶</a></h3><p>Syntax:</p><pre tabindex=0><code>#acc.declare< DataClauseAttr, # dataClause bool # implicit > </code></pre><h4 id=parameters-1>Parameters: <a class=headline-hash href=#parameters-1>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>dataClause</td><td style=text-align:center><code>DataClauseAttr</code></td><td></td></tr><tr><td style=text-align:center>implicit</td><td style=text-align:center><code>bool</code></td><td></td></tr></tbody></table><h3 id=clausedefaultvalueattr>ClauseDefaultValueAttr <a class=headline-hash href=#clausedefaultvalueattr>¶</a></h3><p>DefaultValue Clause</p><p>Syntax:</p><pre tabindex=0><code>#acc.defaultvalue< ::mlir::acc::ClauseDefaultValue # value > </code></pre><p>Enum cases:</p><ul><li>present (<code>Present</code>)</li><li>none (<code>None</code>)</li></ul><h4 id=parameters-2>Parameters: <a class=headline-hash href=#parameters-2>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::ClauseDefaultValue</code></td><td>an enum of type ClauseDefaultValue</td></tr></tbody></table><h3 id=combinedconstructstypeattr>CombinedConstructsTypeAttr <a class=headline-hash href=#combinedconstructstypeattr>¶</a></h3><p>Differentiate between combined constructs</p><p>Syntax:</p><pre tabindex=0><code>#acc.combined_constructs< ::mlir::acc::CombinedConstructsType # value > </code></pre><p>Enum cases:</p><ul><li>kernels_loop (<code>KernelsLoop</code>)</li><li>parallel_loop (<code>ParallelLoop</code>)</li><li>serial_loop (<code>SerialLoop</code>)</li></ul><h4 id=parameters-3>Parameters: <a class=headline-hash href=#parameters-3>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::CombinedConstructsType</code></td><td>an enum of type CombinedConstructsType</td></tr></tbody></table><h3 id=constructattr>ConstructAttr <a class=headline-hash href=#constructattr>¶</a></h3><p>constructs supported by OpenACC</p><p>Syntax:</p><pre tabindex=0><code>#acc.construct< ::mlir::acc::Construct # value > </code></pre><p>Enum cases:</p><ul><li>acc_construct_parallel (<code>acc_construct_parallel</code>)</li><li>acc_construct_kernels (<code>acc_construct_kernels</code>)</li><li>acc_construct_loop (<code>acc_construct_loop</code>)</li><li>acc_construct_data (<code>acc_construct_data</code>)</li><li>acc_construct_enter_data (<code>acc_construct_enter_data</code>)</li><li>acc_construct_exit_data (<code>acc_construct_exit_data</code>)</li><li>acc_construct_host_data (<code>acc_construct_host_data</code>)</li><li>acc_construct_atomic (<code>acc_construct_atomic</code>)</li><li>acc_construct_declare (<code>acc_construct_declare</code>)</li><li>acc_construct_init (<code>acc_construct_init</code>)</li><li>acc_construct_shutdown (<code>acc_construct_shutdown</code>)</li><li>acc_construct_set (<code>acc_construct_set</code>)</li><li>acc_construct_update (<code>acc_construct_update</code>)</li><li>acc_construct_routine (<code>acc_construct_routine</code>)</li><li>acc_construct_wait (<code>acc_construct_wait</code>)</li><li>acc_construct_runtime_api (<code>acc_construct_runtime_api</code>)</li><li>acc_construct_serial (<code>acc_construct_serial</code>)</li></ul><h4 id=parameters-4>Parameters: <a class=headline-hash href=#parameters-4>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::Construct</code></td><td>an enum of type Construct</td></tr></tbody></table><h3 id=dataclauseattr>DataClauseAttr <a class=headline-hash href=#dataclauseattr>¶</a></h3><p>data clauses supported by OpenACC</p><p>Syntax:</p><pre tabindex=0><code>#acc.data_clause< ::mlir::acc::DataClause # value > </code></pre><p>Enum cases:</p><ul><li>acc_copyin (<code>acc_copyin</code>)</li><li>acc_copyin_readonly (<code>acc_copyin_readonly</code>)</li><li>acc_copy (<code>acc_copy</code>)</li><li>acc_copyout (<code>acc_copyout</code>)</li><li>acc_copyout_zero (<code>acc_copyout_zero</code>)</li><li>acc_present (<code>acc_present</code>)</li><li>acc_create (<code>acc_create</code>)</li><li>acc_create_zero (<code>acc_create_zero</code>)</li><li>acc_delete (<code>acc_delete</code>)</li><li>acc_attach (<code>acc_attach</code>)</li><li>acc_detach (<code>acc_detach</code>)</li><li>acc_no_create (<code>acc_no_create</code>)</li><li>acc_private (<code>acc_private</code>)</li><li>acc_firstprivate (<code>acc_firstprivate</code>)</li><li>acc_deviceptr (<code>acc_deviceptr</code>)</li><li>acc_getdeviceptr (<code>acc_getdeviceptr</code>)</li><li>acc_update_host (<code>acc_update_host</code>)</li><li>acc_update_self (<code>acc_update_self</code>)</li><li>acc_update_device (<code>acc_update_device</code>)</li><li>acc_use_device (<code>acc_use_device</code>)</li><li>acc_reduction (<code>acc_reduction</code>)</li><li>acc_declare_device_resident (<code>acc_declare_device_resident</code>)</li><li>acc_declare_link (<code>acc_declare_link</code>)</li><li>acc_cache (<code>acc_cache</code>)</li><li>acc_cache_readonly (<code>acc_cache_readonly</code>)</li></ul><h4 id=parameters-5>Parameters: <a class=headline-hash href=#parameters-5>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::DataClause</code></td><td>an enum of type DataClause</td></tr></tbody></table><h3 id=devicetypeattr>DeviceTypeAttr <a class=headline-hash href=#devicetypeattr>¶</a></h3><p>built-in device type supported by OpenACC</p><p>Syntax:</p><pre tabindex=0><code>#acc.device_type< ::mlir::acc::DeviceType # value > </code></pre><p>Enum cases:</p><ul><li>none (<code>None</code>)</li><li>star (<code>Star</code>)</li><li>default (<code>Default</code>)</li><li>host (<code>Host</code>)</li><li>multicore (<code>Multicore</code>)</li><li>nvidia (<code>Nvidia</code>)</li><li>radeon (<code>Radeon</code>)</li></ul><h4 id=parameters-6>Parameters: <a class=headline-hash href=#parameters-6>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::DeviceType</code></td><td>an enum of type DeviceType</td></tr></tbody></table><h3 id=gangargtypeattr>GangArgTypeAttr <a class=headline-hash href=#gangargtypeattr>¶</a></h3><p>Differentiate the different gang arg values</p><p>Syntax:</p><pre tabindex=0><code>#acc.gang_arg_type< ::mlir::acc::GangArgType # value > </code></pre><p>Enum cases:</p><ul><li>Num (<code>Num</code>)</li><li>Dim (<code>Dim</code>)</li><li>Static (<code>Static</code>)</li></ul><h4 id=parameters-7>Parameters: <a class=headline-hash href=#parameters-7>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::GangArgType</code></td><td>an enum of type GangArgType</td></tr></tbody></table><h3 id=reductionoperatorattr>ReductionOperatorAttr <a class=headline-hash href=#reductionoperatorattr>¶</a></h3><p>built-in reduction operations supported by OpenACC</p><p>Syntax:</p><pre tabindex=0><code>#acc.reduction_operator< ::mlir::acc::ReductionOperator # value > </code></pre><p>Enum cases:</p><ul><li>add (<code>AccAdd</code>)</li><li>mul (<code>AccMul</code>)</li><li>max (<code>AccMax</code>)</li><li>min (<code>AccMin</code>)</li><li>iand (<code>AccIand</code>)</li><li>ior (<code>AccIor</code>)</li><li>xor (<code>AccXor</code>)</li><li>eqv (<code>AccEqv</code>)</li><li>neqv (<code>AccNeqv</code>)</li><li>land (<code>AccLand</code>)</li><li>lor (<code>AccLor</code>)</li></ul><h4 id=parameters-8>Parameters: <a class=headline-hash href=#parameters-8>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>value</td><td style=text-align:center><code>::mlir::acc::ReductionOperator</code></td><td>an enum of type ReductionOperator</td></tr></tbody></table><h3 id=routineinfoattr>RoutineInfoAttr <a class=headline-hash href=#routineinfoattr>¶</a></h3><p>Keeps track of associated acc routine information</p><p>Syntax:</p><pre tabindex=0><code>#acc.routine_info< ::llvm::ArrayRef<SymbolRefAttr> # accRoutines > </code></pre><p>This attribute is used to create the association between a function and its <code>acc.routine</code> operation. A <code>func.func</code> uses this if its name was referenced in an <code>acc routine</code> directive.</p><h4 id=parameters-9>Parameters: <a class=headline-hash href=#parameters-9>¶</a></h4><table><thead><tr><th style=text-align:center>Parameter</th><th style=text-align:center>C++ type</th><th>Description</th></tr></thead><tbody><tr><td style=text-align:center>accRoutines</td><td style=text-align:center><code>::llvm::ArrayRef<SymbolRefAttr></code></td><td></td></tr></tbody></table><h2 id=types-1>Types <a class=headline-hash href=#types-1>¶</a></h2><h3 id=databoundstype>DataBoundsType <a class=headline-hash href=#databoundstype>¶</a></h3><p>Type for representing acc data clause bounds information</p><p>Syntax: <code>!acc.data_bounds_ty</code></p><h3 id=declaretokentype>DeclareTokenType <a class=headline-hash href=#declaretokentype>¶</a></h3><p>declare token type</p><p>Syntax: <code>!acc.declare_token</code></p><p><code>acc.declare_token</code> is a type returned by a <code>declare_enter</code> operation and can be passed to a <code>declare_exit</code> operation to represent an implicit data region.</p><h2 id=enums>Enums <a class=headline-hash href=#enums>¶</a></h2><h3 id=clausedefaultvalue>ClauseDefaultValue <a class=headline-hash href=#clausedefaultvalue>¶</a></h3><p>DefaultValue Clause</p><h4 id=cases>Cases: <a class=headline-hash href=#cases>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>Present</td><td style=text-align:center><code>0</code></td><td>present</td></tr><tr><td style=text-align:center>None</td><td style=text-align:center><code>1</code></td><td>none</td></tr></tbody></table><h3 id=combinedconstructstype>CombinedConstructsType <a class=headline-hash href=#combinedconstructstype>¶</a></h3><p>Differentiate between combined constructs</p><h4 id=cases-1>Cases: <a class=headline-hash href=#cases-1>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>KernelsLoop</td><td style=text-align:center><code>1</code></td><td>kernels_loop</td></tr><tr><td style=text-align:center>ParallelLoop</td><td style=text-align:center><code>2</code></td><td>parallel_loop</td></tr><tr><td style=text-align:center>SerialLoop</td><td style=text-align:center><code>3</code></td><td>serial_loop</td></tr></tbody></table><h3 id=construct>Construct <a class=headline-hash href=#construct>¶</a></h3><p>constructs supported by OpenACC</p><h4 id=cases-2>Cases: <a class=headline-hash href=#cases-2>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>acc_construct_parallel</td><td style=text-align:center><code>0</code></td><td>acc_construct_parallel</td></tr><tr><td style=text-align:center>acc_construct_kernels</td><td style=text-align:center><code>1</code></td><td>acc_construct_kernels</td></tr><tr><td style=text-align:center>acc_construct_loop</td><td style=text-align:center><code>2</code></td><td>acc_construct_loop</td></tr><tr><td style=text-align:center>acc_construct_data</td><td style=text-align:center><code>3</code></td><td>acc_construct_data</td></tr><tr><td style=text-align:center>acc_construct_enter_data</td><td style=text-align:center><code>4</code></td><td>acc_construct_enter_data</td></tr><tr><td style=text-align:center>acc_construct_exit_data</td><td style=text-align:center><code>5</code></td><td>acc_construct_exit_data</td></tr><tr><td style=text-align:center>acc_construct_host_data</td><td style=text-align:center><code>6</code></td><td>acc_construct_host_data</td></tr><tr><td style=text-align:center>acc_construct_atomic</td><td style=text-align:center><code>7</code></td><td>acc_construct_atomic</td></tr><tr><td style=text-align:center>acc_construct_declare</td><td style=text-align:center><code>8</code></td><td>acc_construct_declare</td></tr><tr><td style=text-align:center>acc_construct_init</td><td style=text-align:center><code>9</code></td><td>acc_construct_init</td></tr><tr><td style=text-align:center>acc_construct_shutdown</td><td style=text-align:center><code>10</code></td><td>acc_construct_shutdown</td></tr><tr><td style=text-align:center>acc_construct_set</td><td style=text-align:center><code>11</code></td><td>acc_construct_set</td></tr><tr><td style=text-align:center>acc_construct_update</td><td style=text-align:center><code>12</code></td><td>acc_construct_update</td></tr><tr><td style=text-align:center>acc_construct_routine</td><td style=text-align:center><code>13</code></td><td>acc_construct_routine</td></tr><tr><td style=text-align:center>acc_construct_wait</td><td style=text-align:center><code>14</code></td><td>acc_construct_wait</td></tr><tr><td style=text-align:center>acc_construct_runtime_api</td><td style=text-align:center><code>15</code></td><td>acc_construct_runtime_api</td></tr><tr><td style=text-align:center>acc_construct_serial</td><td style=text-align:center><code>16</code></td><td>acc_construct_serial</td></tr></tbody></table><h3 id=dataclause>DataClause <a class=headline-hash href=#dataclause>¶</a></h3><p>data clauses supported by OpenACC</p><h4 id=cases-3>Cases: <a class=headline-hash href=#cases-3>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>acc_copyin</td><td style=text-align:center><code>1</code></td><td>acc_copyin</td></tr><tr><td style=text-align:center>acc_copyin_readonly</td><td style=text-align:center><code>2</code></td><td>acc_copyin_readonly</td></tr><tr><td style=text-align:center>acc_copy</td><td style=text-align:center><code>3</code></td><td>acc_copy</td></tr><tr><td style=text-align:center>acc_copyout</td><td style=text-align:center><code>4</code></td><td>acc_copyout</td></tr><tr><td style=text-align:center>acc_copyout_zero</td><td style=text-align:center><code>5</code></td><td>acc_copyout_zero</td></tr><tr><td style=text-align:center>acc_present</td><td style=text-align:center><code>6</code></td><td>acc_present</td></tr><tr><td style=text-align:center>acc_create</td><td style=text-align:center><code>7</code></td><td>acc_create</td></tr><tr><td style=text-align:center>acc_create_zero</td><td style=text-align:center><code>8</code></td><td>acc_create_zero</td></tr><tr><td style=text-align:center>acc_delete</td><td style=text-align:center><code>9</code></td><td>acc_delete</td></tr><tr><td style=text-align:center>acc_attach</td><td style=text-align:center><code>10</code></td><td>acc_attach</td></tr><tr><td style=text-align:center>acc_detach</td><td style=text-align:center><code>11</code></td><td>acc_detach</td></tr><tr><td style=text-align:center>acc_no_create</td><td style=text-align:center><code>12</code></td><td>acc_no_create</td></tr><tr><td style=text-align:center>acc_private</td><td style=text-align:center><code>13</code></td><td>acc_private</td></tr><tr><td style=text-align:center>acc_firstprivate</td><td style=text-align:center><code>14</code></td><td>acc_firstprivate</td></tr><tr><td style=text-align:center>acc_deviceptr</td><td style=text-align:center><code>15</code></td><td>acc_deviceptr</td></tr><tr><td style=text-align:center>acc_getdeviceptr</td><td style=text-align:center><code>16</code></td><td>acc_getdeviceptr</td></tr><tr><td style=text-align:center>acc_update_host</td><td style=text-align:center><code>17</code></td><td>acc_update_host</td></tr><tr><td style=text-align:center>acc_update_self</td><td style=text-align:center><code>18</code></td><td>acc_update_self</td></tr><tr><td style=text-align:center>acc_update_device</td><td style=text-align:center><code>19</code></td><td>acc_update_device</td></tr><tr><td style=text-align:center>acc_use_device</td><td style=text-align:center><code>20</code></td><td>acc_use_device</td></tr><tr><td style=text-align:center>acc_reduction</td><td style=text-align:center><code>21</code></td><td>acc_reduction</td></tr><tr><td style=text-align:center>acc_declare_device_resident</td><td style=text-align:center><code>22</code></td><td>acc_declare_device_resident</td></tr><tr><td style=text-align:center>acc_declare_link</td><td style=text-align:center><code>23</code></td><td>acc_declare_link</td></tr><tr><td style=text-align:center>acc_cache</td><td style=text-align:center><code>24</code></td><td>acc_cache</td></tr><tr><td style=text-align:center>acc_cache_readonly</td><td style=text-align:center><code>25</code></td><td>acc_cache_readonly</td></tr></tbody></table><h3 id=devicetype>DeviceType <a class=headline-hash href=#devicetype>¶</a></h3><p>built-in device type supported by OpenACC</p><h4 id=cases-4>Cases: <a class=headline-hash href=#cases-4>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>None</td><td style=text-align:center><code>0</code></td><td>none</td></tr><tr><td style=text-align:center>Star</td><td style=text-align:center><code>1</code></td><td>star</td></tr><tr><td style=text-align:center>Default</td><td style=text-align:center><code>2</code></td><td>default</td></tr><tr><td style=text-align:center>Host</td><td style=text-align:center><code>3</code></td><td>host</td></tr><tr><td style=text-align:center>Multicore</td><td style=text-align:center><code>4</code></td><td>multicore</td></tr><tr><td style=text-align:center>Nvidia</td><td style=text-align:center><code>5</code></td><td>nvidia</td></tr><tr><td style=text-align:center>Radeon</td><td style=text-align:center><code>6</code></td><td>radeon</td></tr></tbody></table><h3 id=gangargtype>GangArgType <a class=headline-hash href=#gangargtype>¶</a></h3><p>Differentiate the different gang arg values</p><h4 id=cases-5>Cases: <a class=headline-hash href=#cases-5>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>Num</td><td style=text-align:center><code>0</code></td><td>Num</td></tr><tr><td style=text-align:center>Dim</td><td style=text-align:center><code>1</code></td><td>Dim</td></tr><tr><td style=text-align:center>Static</td><td style=text-align:center><code>2</code></td><td>Static</td></tr></tbody></table><h3 id=reductionoperator>ReductionOperator <a class=headline-hash href=#reductionoperator>¶</a></h3><p>built-in reduction operations supported by OpenACC</p><h4 id=cases-6>Cases: <a class=headline-hash href=#cases-6>¶</a></h4><table><thead><tr><th style=text-align:center>Symbol</th><th style=text-align:center>Value</th><th>String</th></tr></thead><tbody><tr><td style=text-align:center>AccAdd</td><td style=text-align:center><code>0</code></td><td>add</td></tr><tr><td style=text-align:center>AccMul</td><td style=text-align:center><code>1</code></td><td>mul</td></tr><tr><td style=text-align:center>AccMax</td><td style=text-align:center><code>2</code></td><td>max</td></tr><tr><td style=text-align:center>AccMin</td><td style=text-align:center><code>3</code></td><td>min</td></tr><tr><td style=text-align:center>AccIand</td><td style=text-align:center><code>4</code></td><td>iand</td></tr><tr><td style=text-align:center>AccIor</td><td style=text-align:center><code>5</code></td><td>ior</td></tr><tr><td style=text-align:center>AccXor</td><td style=text-align:center><code>6</code></td><td>xor</td></tr><tr><td style=text-align:center>AccEqv</td><td style=text-align:center><code>7</code></td><td>eqv</td></tr><tr><td style=text-align:center>AccNeqv</td><td style=text-align:center><code>8</code></td><td>neqv</td></tr><tr><td style=text-align:center>AccLand</td><td style=text-align:center><code>9</code></td><td>land</td></tr><tr><td style=text-align:center>AccLor</td><td style=text-align:center><code>10</code></td><td>lor</td></tr></tbody></table><div class=edit-meta><br></div><nav class=pagination><a class="nav nav-prev" href=https://mlir.llvm.org/docs/Dialects/DLTITransformOps/ title><i class="fas fa-arrow-left" aria-hidden=true></i> Prev -</a> <a class="nav nav-next" href=https://mlir.llvm.org/docs/Dialects/Affine/ title="'affine' Dialect">Next - 'affine' Dialect <i class="fas fa-arrow-right" aria-hidden=true></i></a></nav><footer><p class=powered>Powered by <a href=https://gohugo.io>Hugo</a>. Theme by <a href=https://themes.gohugo.io/hugo-theme-techdoc/>TechDoc</a>. Designed by <a href=https://github.com/thingsym/hugo-theme-techdoc>Thingsym</a>.</p></footer></main><div class=sidebar><nav class=slide-menu><ul><li><a href=https://mlir.llvm.org/>Home</a></li><li><a href=https://mlir.llvm.org/users/>Users of MLIR</a></li><li><a href=https://mlir.llvm.org/pubs/>MLIR Related Publications</a></li><li><a href=https://mlir.llvm.org/talks/>Talks</a></li><li><a href=https://mlir.llvm.org/deprecation/>Deprecations & Current Refactoring</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/getting_started/>Getting Started<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/getting_started/ReportingIssues/>Reporting Issues</a></li><li><a href=https://mlir.llvm.org/getting_started/Debugging/>Debugging Tips</a></li><li><a href=https://mlir.llvm.org/getting_started/Faq/>FAQ</a></li><li><a href=https://mlir.llvm.org/getting_started/Contributing/>How to Contribute</a></li><li><a href=https://mlir.llvm.org/getting_started/DeveloperGuide/>Developer Guide</a></li><li><a href=https://mlir.llvm.org/getting_started/openprojects/>Open Projects</a></li><li><a href=https://mlir.llvm.org/getting_started/Glossary/>Glossary</a></li><li><a href=https://mlir.llvm.org/getting_started/TestingGuide/>Testing Guide</a></li></ul></li><li class="parent has-sub-menu"><a href=https://mlir.llvm.org/docs/>Code Documentation<span class="mark opened">-</span></a><ul class=sub-menu><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Bindings/>Bindings<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Bindings/Python/>MLIR Python Bindings</a></li></ul></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Tools/>Tools<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Tools/MLIRLSP/>MLIR : Language Server Protocol</a></li><li><a href=https://mlir.llvm.org/docs/Tools/mlir-reduce/>MLIR Reduce</a></li><li><a href=https://mlir.llvm.org/docs/Tools/mlir-rewrite/>mlir-rewrite</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/QuantPasses/></a></li><li><a href=https://mlir.llvm.org/docs/ActionTracing/>Action: Tracing and Debugging MLIR-based Compilers</a></li><li><a href=https://mlir.llvm.org/docs/BufferDeallocationInternals/>Buffer Deallocation - Internals</a></li><li><a href=https://mlir.llvm.org/docs/Bufferization/>Bufferization</a></li><li><a href=https://mlir.llvm.org/docs/DataLayout/>Data Layout Modeling</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/DefiningDialects/>Defining Dialects<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/DefiningDialects/Constraints/>Constraints</a></li><li><a href=https://mlir.llvm.org/docs/DefiningDialects/AttributesAndTypes/>Defining Dialect Attributes and Types</a></li><li><a href=https://mlir.llvm.org/docs/DefiningDialects/Operations/>Operation Definition Specification (ODS)</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/Diagnostics/>Diagnostic Infrastructure</a></li><li><a href=https://mlir.llvm.org/docs/DialectConversion/>Dialect Conversion</a></li><li class="parent has-sub-menu"><a href=https://mlir.llvm.org/docs/Dialects/>Dialects<span class="mark opened">-</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Dialects/DLTITransformOps/></a></li><li class=active><a href=https://mlir.llvm.org/docs/Dialects/OpenACCDialect/>'acc' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Affine/>'affine' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/AMDGPU/>'amdgpu' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/AMX/>'amx' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ArithOps/>'arith' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ArmNeon/>'arm_neon' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ArmSVE/>'arm_sve' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ArmSME/>'ArmSME' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/AsyncDialect/>'async' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/BufferizationOps/>'bufferization' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ControlFlowDialect/>'cf' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ComplexOps/>'complex' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/DLTIDialect/>'dlti' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/EmitC/>'emitc' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Func/>'func' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/GPU/>'gpu' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/IndexOps/>'index' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/IRDL/>'irdl' Dialect</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Dialects/Linalg/>'linalg' Dialect<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Dialects/Linalg/OpDSL/>Linalg OpDSL</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/Dialects/LLVM/>'llvm' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/MathOps/>'math' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/MemRef/>'memref' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Mesh/>'mesh' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/MLProgramOps/>'ml_program' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/MPI/>'mpi' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/NVGPU/>'nvgpu' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/NVVMDialect/>'nvvm' Dialect</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Dialects/OpenMPDialect/>'omp' Dialect<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Dialects/OpenMPDialect/ODS/>ODS Documentation</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/Dialects/PDLInterpOps/>'pdl_interp' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/PDLOps/>'pdl' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/PolynomialDialect/>'polynomial' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/PtrOps/>'ptr' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/QuantDialect/>'quant' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ROCDLDialect/>'rocdl' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/SCFDialect/>'scf' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/ShapeDialect/>'shape' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/SparseTensorOps/>'sparse_tensor' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/TensorOps/>'tensor' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/UBOps/>'ub' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/VCIXDialect/>'vcix' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Vector/>'vector' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/X86Vector/>'x86vector' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/XeGPU/>'xegpu' Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Builtin/>Builtin Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/MatchOpInterfaces/>OpInterface definitions</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/SPIR-V/>SPIR-V Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/TOSA/>Tensor Operator Set Architecture (TOSA) Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Dialects/Transform/>Transform Dialect</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/Interfaces/>Interfaces</a></li><li><a href=https://mlir.llvm.org/docs/TargetLLVMIR/>LLVM IR Target</a></li><li><a href=https://mlir.llvm.org/docs/BytecodeFormat/>MLIR Bytecode Format</a></li><li><a href=https://mlir.llvm.org/docs/CAPI/>MLIR C API</a></li><li><a href=https://mlir.llvm.org/docs/LangRef/>MLIR Language Reference</a></li><li><a href=https://mlir.llvm.org/docs/ReleaseNotes/>MLIR Release Notes</a></li><li><a href=https://mlir.llvm.org/docs/Canonicalization/>Operation Canonicalization</a></li><li><a href=https://mlir.llvm.org/docs/OwnershipBasedBufferDeallocation/>Ownership-based Buffer Deallocation</a></li><li><a href=https://mlir.llvm.org/docs/PassManagement/>Pass Infrastructure</a></li><li><a href=https://mlir.llvm.org/docs/Passes/>Passes</a></li><li><a href=https://mlir.llvm.org/docs/PatternRewriter/>Pattern Rewriting : Generic DAG-to-DAG Rewriting</a></li><li><a href=https://mlir.llvm.org/docs/PDLL/>PDLL - PDL Language</a></li><li><a href=https://mlir.llvm.org/docs/Quantization/>Quantization</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Rationale/>Rationale<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Rationale/RationaleGenericDAGRewriter/>Generic DAG Rewriter Infrastructure Rationale</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/RationaleLinalgDialect/>Linalg Dialect Rationale: The Case For Compiler-Friendly Custom Operations</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/Rationale/>MLIR Rationale</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/MLIRForGraphAlgorithms/>MLIR: Incremental Application to Graph Algorithms in ML Frameworks</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/RationaleSimplifiedPolyhedralForm/>MLIR: The case for a simplified polyhedral form</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/SideEffectsAndSpeculation/>Side Effects & Speculation</a></li><li><a href=https://mlir.llvm.org/docs/Rationale/UsageOfConst/>Usage of 'const' in MLIR, for core IR types</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/ShapeInference/>Shape Inference</a></li><li><a href=https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/>SPIR-V Dialect to LLVM Dialect conversion manual</a></li><li><a href=https://mlir.llvm.org/docs/SymbolsAndSymbolTables/>Symbols and Symbol Tables</a></li><li><a href=https://mlir.llvm.org/docs/DeclarativeRewrites/>Table-driven Declarative Rewrite Rule (DRR)</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Traits/>Traits<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Traits/Broadcastable/>The `Broadcastable` Trait</a></li></ul></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Tutorials/>Tutorials<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Tutorials/CreatingADialect/>Creating a Dialect</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/QuickstartRewrites/>Quickstart tutorial to adding MLIR graph rewrite</a></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Tutorials/Toy/>Toy Tutorial<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-1/>Chapter 1: Toy Language and AST</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-2/>Chapter 2: Emitting Basic MLIR</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-3/>Chapter 3: High-level Language-Specific Analysis and Transformation</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-4/>Chapter 4: Enabling Generic Transformation with Interfaces</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-5/>Chapter 5: Partial Lowering to Lower-Level Dialects for Optimization</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-6/>Chapter 6: Lowering to LLVM and CodeGeneration</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/Toy/Ch-7/>Chapter 7: Adding a Composite Type to Toy</a></li></ul></li><li class=has-sub-menu><a href=https://mlir.llvm.org/docs/Tutorials/transform/>Transform Dialect Tutorial<span class="mark closed">+</span></a><ul class=sub-menu><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/Ch0/>Chapter 0: A Primer on “Structured” Linalg Operations</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/Ch1/>Chapter 1: Combining Existing Transformations</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/Ch2/>Chapter 2: Adding a Simple New Transformation Operation</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/Ch3/>Chapter 3: More than Simple Transform Operations</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/Ch4/>Chapter 4: Matching Payload with Transform Operations</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/transform/ChH/>Chapter H: Reproducing Halide Schedule</a></li></ul></li><li><a href=https://mlir.llvm.org/docs/Tutorials/UnderstandingTheIRStructure/>Understanding the IR Structure</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/MlirOpt/>Using `mlir-opt`</a></li><li><a href=https://mlir.llvm.org/docs/Tutorials/DataFlowAnalysis/>Writing DataFlow Analyses in MLIR</a></li></ul></li></ul></li></ul></nav><div class=sidebar-footer></div></div></div><a href=# id=backtothetop-fixed class=backtothetop data-backtothetop-duration=600 data-backtothetop-easing=easeOutQuart data-backtothetop-fixed-fadein=1000 data-backtothetop-fixed-fadeout=1000 data-backtothetop-fixed-bottom=10 data-backtothetop-fixed-right=20><span class="fa-layers fa-fw"><i class="fas fa-circle"></i> <i class="fas fa-arrow-circle-up"></i></span></a></div></body></html>