Revert #196094 and its follow-up #209701. The expression override
mechanism does not preserve reduction-object identity across all data
environments, leading to incorrect lowering for procedure-local and
nested private arrays.
Restore the pre-#196094 lowering. Keep coverage showing that
array-element constructs compile through the boxed-array reduction
path.
This was in response to this comment:
https://github.com/llvm/llvm-project/pull/196094#issuecomment-5248085956
Fixing exactly the bug in the comment wasn't hard but AI code review
found a large number of follow on bugs so I think the design needs a
rethink, and definitely shouldn't be included in the LLVM release.
Assisted-by: Codex
The TODO message in `verify()` was not completly displayed, so it has
been replaced by `emitErrorOp`.
In addition, the test in `ConvertVariable` has been updated by adding
pointer direct component case.
Fix#193829
(cherry picked from commit 9adc2537177d0813301b3eab1bae906735c984a4)
This PR resolves the issue #205953. Its purpose is to correct the use of
`team_type` by lowering it to Addr rather than to Box.
(cherry picked from commit 0563f0d88d8314edcde7deab465759a84696e894)
Mark eligible `INTENT(IN)` dummy data objects with a `fir.read_only`
argument attribute during lowering, and teach `FunctionAttr` to
translate the marker to LLVM `readonly` for reference arguments when
optimizing for speed.
LLVM FunctionAttrs can infer `readonly` later in the pipeline, but this
happens too late for IPSCCP function specialization.
`FunctionSpecializer::promoteConstantStackValues` requires
`onlyReadsMemory()` before promoting constant by-reference scalar
arguments. Providing `readonly` satisfies this check and allows these
arguments to be promoted, enabling function specialization, including in
non-LTO builds.
For direct-data arguments, restrict propagation to by-reference
non-character intrinsic scalars. Ordinary arrays and derived-type
arguments are excluded because compiler-generated copy-out may write
back through a forwarded dummy argument. LLVM `readonly` does not permit
such writes, even when the stored value is unchanged.
`INTENT(IN)` `POINTER` and `ALLOCATABLE` dummies are supported with
shallow descriptor semantics: `readonly` protects the descriptor
storage, but does not imply that data addressed by the descriptor is
read-only. In particular, a `POINTER` target may still be modified.
Keep `readonly` handling independent from `noalias` and `nocapture`.
`TARGET` arguments may receive `readonly`, but must not receive
`noalias` or `nocapture`. Apply `readonly` to `BIND(C)` definitions
while leaving external C declarations unannotated. `VALUE`,
`ASYNCHRONOUS`, and `VOLATILE` dummy arguments remain excluded.
In our measurements, the enabled function specialization gives a 1.5x
speedup on 548.exchange2 on NVIDIA Neoverse V2 in a non-LTO
configuration.
Update existing lowering tests to account for the new `fir.read_only`
argument attribute, and add focused coverage:
* `dummy-argument-readonly.f90` checks marker emission for supported
scalar and descriptor arguments, excluded cases, and shallow pointer
semantics.
* `function-attrs-readonly.fir` checks translation from `fir.read_only`
to LLVM `readonly`, including `TARGET`, `BIND(C)`, and descriptor
references.
* `function-attr-readonly.f90` checks the optimization-level pipeline
gate and verifies that `readonly` reaches final LLVM IR.
* `call-copy-in-out.f90` checks that ordinary array copy-out cases
remain unmarked and that pointer-target copy-out is compatible with
shallow descriptor `readonly`.
In https://github.com/llvm/llvm-project/pull/206525 the runtime was set
to default alloctable and pointer arrays to 64byte alignment. In this commit
the inlineallocation case is handled.
The same approach was taken and an alignment attribute was added,
with logic to promote to aligned_alloc or posixMemalign for apple
targets, when requested alignment is greater than 16.
Of note alignof(std::max_align_t)) was not used as the runtime target
may be different than the compile target, so 16 was hardcoded as when to
use aligned_alloc
Assisted-by: Opus 4.8
Add lowering support for OpenMP DECLARE VARIANT. Each directive is
recorded on its base procedure during the structure checks. During call
lowering, best variant is picked based on recorded selecter and context.
Note: this PR does not include module-file support for DECLARE VARIANT,
so cross-module variant resolution is not yet available; it will be
added in a follow-up PR.
Assisted by: Cursor
The goal of this PR is to add support for allocating NON-allocatable
coarrays that are not in the `Main` program.
All allocations will be performed in a function called
`__mif_save_coarrays_allocate`. If this function exists, then
LLVM::GlobalCtors is created (or modified).
The initialization of the MIF with `mif.init` is moved too in a
GlobalCtors with a priority of 0 to ensure that the initialization takes
place before any MIF allocations.
---------
Co-authored-by: Dan Bonachea <dobonachea@lbl.gov>
FIR is extremely verbose with derived types because the type lists all
the component names and types, so in case of nested derived type, the
type length in the assembly may grow over thousands of characters.
This patch implements the `getAlias` that allows using the fir.type name
as the MLIR type alias. This significantly reduce the IR size in
application with derived type.
I am planning to enable this by default, but there will be some churn in
the lit test, so I am first adding it under a developer option.
Assisted-by: Cursor
Use the source array extent when advancing cumulative strides across
scalar subscripts, whose selected extent is one. Add coverage for mixed
scalar and triplet dimensions.
Add an opt-in cache to fir::AliasAnalysis that memoizes getSource()
results keyed on (value, flags). Caching is off by default and turned on
via enableSourceCache(); getSource() becomes a thin wrapper over the
uncached getSourceImpl(). The cache is a frozen snapshot with no
automatic invalidation and lives no longer than the AliasAnalysis
instance, so a client enables it only for a region in which it does not
mutate IR in a way that would change a source.
Flang's LICM enables the cache on its fir::AliasAnalysis before adding
it to the mlir::AliasAnalysis aggregate: LICM only moves operations, so
getSource()'s inputs are unchanged across the hoists, and the pass
manager drops the analysis (and its cache) after the pass.
Add a unit test covering the default-off behavior, cache fill/hit, and
that disableSourceCache() clears and bypasses the cache.
Assisted-by: Cursor
Add an attribute which indicates that a call is to an intrinsic
function. Intrinsic here means that the function declaration was found
in a module found in the intrinsic path.
This allows us to, when optimizing, make sure that we correctly identify
calls to intrinsic functions whose behaviours we know, and not
user-defined functions whose names happen to match.
Inlining can leave a `fir.declare`'s dummy_scope operand (`!fir.dscope`)
defined outside an OpenACC compute region while the declare itself gets
rematerialized inside, since `fir.dummy_scope` wasn't registered as an
OutlineRematerializationOpInterface candidate.
This dangling, unmappable live-in is illegal for parallelization.
Fix: Added `OutlineIdentityOperandOpInterface`, implemented by
`fir.declare/hlfir.declare`, so `OffloadLiveInValueCanonicalization`
drops their dummy_scope operand (an identity token that must never be
duplicated) instead of cloning it when sinking/rematerializing them into
offload regions.
This patch adds a parser and basic checks for the Apply clause from
OpenMP 6.0 in Flang.
It is missing semantic checks for nesting loop transformations that are
being worked on.
Move towards using the triple for representing incompatible
ISA changes. Use the subarch field to represent the various
incompatible cases. Previously we pretended a single triple arch
was universally compatible, and only distinguished by function
level subtargets. Move towards using distinct triples to enable
more sophisticated toolchain handling in the future, like proper
runtime library linking.
Introduce a new subarch per unique ISA, but also introduce
"major subarches" which are compatible by a set of covered
minor ISA versions. These map to the existing generic targets.
There are a few placeholder subarch entries, which currently
have missing backing generic arches for codegen.
This should be the preferred triple arch name going forward,
but is treated as an alias of amdgcn. This does not yet change
clang to emit the new triples.
Part of #154925
The MLIR pass that removes operations exclusively intended for the host
from OpenMP target offload modules is currently defined as part of
Flang. However, this is a feature that would benefit from being reusable
by other frontends, as removing such operations is a requirement for all
OpenMP target device modules prior to LLVM IR translation.
By moving the `omp-host-op-filtering` pass out of Flang, it had to be
updated to work on a lower-level LLVM dialect-based representation,
rather than FIR. This simplified some of the existing edge cases, such
as `fir.declare` ops and `fir.boxchar` type handling. In addition, new
function arguments are introduced as placeholders and return values from
host-only functions are removed, producing a cleaner result and
simplifying the pass as compared to previously.
As a result of a later execution of this pass, dynamic dispatch of host
functions via dispatch table using `fir.dispatch`, `fir.type_info` and
`fir.dt_entry` ops would break due to the removal of `fir.dt_entry`
operations pointing to deleted host functions, while `fir.dispatch` ops
pointing to the same functions would remain. The
`omp-function-filtering` pass is updated to prevent `fir.dt_entry` from
being deleted and let them point to an undefined symbol. This lets FIR
lowering to LLVM work well enough to get to the host op filtering stage,
where any resulting operations for this host dynamic dispatch are
cleaned up.
It is important we recognize acc_on_device calls as they need to be
folded during compilation. Emitting this operation helps with the
recognition of the runtime call in the optimizer.
```fortran
module m
integer, device :: aaa; bind(c, name='aaa_from_c') :: aaa
integer, constant :: zzz; bind(c, name='zzz_from_c') :: zzz
end module
```
```c
extern int aaa_from_c;
#pragma acc declare create(aaa_from_c) // same for zzz_from_c
```
In this code, under `-gpu=mem:unified` a module-scope device/constant
global is registered only as a CUDA variable. Another translation unit
that declares the same symbol on the host can make it be treated as host
memory, so the device symbol is unresolved at run time.
Fix: emit `cuf.register_variable_static` for device/constant globals
under `mem:unified` so they are additionally registered as
device-resident.
Fixes#199092
Flang previously rejected intervening code between associated loops in a
collapsed nest (e.g. `collapse(2)` with statements between the outer and
inner DO). This patch removes that restriction and implements correct
lowering.
**Semantics:** accept intervening code retroactively for all versions,
except when perfect nesting is still required
OpenMP 5.0 introduced support for collapsing imperfectly nested loops
for worksharing-loop, simd, taskloop, and distribute constructs. OpenMP
5.1 later formalized this under the Canonical Loop Nest (CLN)
definition. This support is applied retroactively for all OpenMP
versions, since the semantics are safe to implement regardless of
version and for compatibility with other compilers.
The only case where perfect nesting is still enforced is when ordered
semantics require it:
- Pre-5.2: any ordered clause with an argument requires perfect nesting.
- 5.2+: perfect nesting is required only when the loop body contains
ordered directives with doacross (or the legacy depend(sink/source))
clauses.
A `DoacrossFinder` visitor is added to detect doacross directives in the
loop body while correctly not descending into nested OpenMP block/loop
constructs (which have their own binding context).
**Lowering**
Intervening statements are guarded by induction variable comparisons
within the flat `omp.loop_nest` body: "before" code executes when inner
induction variables equal their lower bounds; "after" code executes when
inner induction variables equal their last iteration values (`lb + ((ub
- lb) / step) * step)` for non-unit steps).
**Testing**
- Lowering lit test added
- Semantics tests: updated to remove expected errors for now-accepted
code; new tests added for coverage.
- llvm-test-suite update: (will add once opened)
Assisted by: Copilot
Add a `defer-acc-routines` option to `cuf-function-rewrite`. When set,
`on_device()` is not folded in the host copy of an OpenACC routine (has
`acc.routine_info`, not in a `gpu.module`), because that body is later
cloned into the device routine and would otherwise bake in the host
value (`.false.`). A later run folds each copy in its own context. Calls
already in a `gpu.module` are still folded.
Currently, when an ArrayElement is used within a Reduction clause, it
will be lowered with the reduction referencing the box containing the
array, not just the element.
To address this, adjust Flang lowering to track expressions alongside
symbol to ensure that the Array Element context is not lost and
considered when lowering a reduction with Array Element. This ensures
that, when represented in HLFIR, it will be just the element's type,
rather than the full array.
Currently this excludes DO CONCURRENT as it excludes Array Elements, and
is limited to Array Elements but there are options to expand this into
Array Sections in the future.
Assisted-by: Codex
This is opt-in by an engineering option and disabled by default.
In section 10.1.5.2.4 of the 2023 Fortran standard "Evaluation of
numerical intrinsic operations", the standard explicitly allows
alternate mathematically equivalent lowerings. For example the source
expression X + Y + Z could be evaluated (X + Y) + Z, X + (Y + Z) or even
(X + Z) + Y, etc.
The open source benchmark SNBone shows significantly better results with
classic flang because classic flang emits real arithmetic expressions in
a different order. In the case of this benchmark it reduces dependency
depth for instructions issued to the vector unit, allowing for more of
the arithmetic to be parallelised over multiple vector execution units
in the ALU.
The lowering added by this patch tries to mimic the way classic flang
orders instructions for these expressions. I did not read any classic
flang source when writing this patch. There is still a notable
difference in that classic flang uses FMA intrinsics whereas LLVM Flang
relies on the rest of the pipeline to introduce FMA when it is safe to
do so.
This is a much less aggressive optimisation than simply enabling reassoc
in the fast-math flags because it does not allow reassociation between
Fortran language statements. This is why I implemented it in lowering.
The new option enables an experimental lowering path for scalar real
top-level addition chains. When enabled with
-enable-split-sum-expression-tree-lowering, eligible sums are split
after the first two terms and rebuilt as a right-associated tail plus
head. This lets the independent tail terms be evaluated before the
assignment-related head, giving the backend a different expression tree
while leaving the default lowering unchanged.
The transform is deliberately narrow. It only applies to scalar real RHS
expressions in assignments and rejects cases with vector subscripts,
parentheses, subtraction, procedure references, or volatile/asynchronous
symbols on either side of the assignment. Subtraction is left out
because the split would need to carry signed terms; division stays
within an individual additive term and does not change the top-level
chain.
In testing I have found some tests in the Fujitsu test suite miscompare
due to small changes in floating point rounding. There are no failures
or regressions in SPEC2017 or SPEC2026. I think it would be legal
according to the Fortran standard to enable this by default, but I am
not proposing that here, and will not consider it until after the LLVM
release branch point.
Assisted-by: Codex
Audit of every #include directive in flang/lib/ and
flang/include/flang/:
- Remove 423 #include directives (293 from lib/ translation units, 23
from lib-private headers, 107 from public headers) that are not needed:
no entity declared by the removed header is used - directly or in any
conditional-compilation branch (Windows, macOS, ARM/PowerPC,
REAL(16)/REAL(2) configuration macros) - in the including file, and for
headers no consumer in flang/lib, flang/tools, flang/unittests, or
flang-rt relies on the include as a re-export. Load-bearing includes
were examined and retained: textual .inc/.def expansions, tablegen
pass-header declaration/definition pairings, dialect headers carrying
generated-op interface dependencies, config-macro providers (e.g.
float128.h ahead of matmul-instances.inc guards), and ADL formatting
machinery (Evaluate/formatting.h).
- Add 5 missing direct #includes in lib translation units that used
entities only via now-removed transitive includes, and drop one
vestigial 'using namespace Fortran::runtime;' whose declaring header no
longer arrives transitively (nothing namespace-qualified from it was
used), per the include rules in flang/docs/C++style.md.
Assisted-by: AI
Add a general acc utility for computing the size and alignment of a
type. It works for simple scalar types as well as nested and aggregate
types like arrays, tuples, and structures.
Because some types come from other dialects, the utility can hand those
off to a specialized helper that understands them. This lets sizing work
seamlessly even for mixed types, such as an aggregate whose members come
from a different dialect.
Add a Fortran-specific helper so Fortran types are sized correctly,
falling back to the general utility for everything else.
Include unit tests covering a range of scenarios, including scalars,
arrays, aggregates, and mixed-dialect types.
Move the DSA helper from the private OmpAttributeVisitor::SetSymbolDSA
into the public openmp-dsa header next to GetSymbolDSA, and share the
DSA flag set through a single GetDataSharingAttributeFlags().
This lets an upcoming metadirective lowering change reuse the helper to
set the predetermined DSA of the loop induction variables of a selected
variant. A loop-associated variant is resolved during lowering, so the
usual semantic DSA resolution never runs on its loop nest and lowering
must set those flags itself.
Assisted with Copilot.
Add parser and semantics plumbing for configured implicit USE modules,
reusing
the normal USE-association path so public symbols and common blocks are
imported
consistently.
Move FlangOMPContext from lowering into semantics as the
MLIR-independent OmpVariantMatchContext. This lets an upcoming semantic
check reuse the matcher to skip the loop-nest checks for a metadirective
variant that cannot be selected on the current target, such as a
`device={kind(nohost)}` variant in a host compilation. No functional
change to metadirective lowering.
Assisted with copilot.
The minimum-blocks-per-multiprocessor (second) and maximum-blocks-per-cluster
(third) operands of the CUDA Fortran `LAUNCH_BOUNDS()` prefix are optional, as in
CUDA C++ `__launch_bounds__`. A single-operand `LAUNCH_BOUNDS(N)` is now accepted;
1, 2, or 3 integer constants are valid.
Previously `LAUNCH_BOUNDS(N)` with a single operand was rejected, and the
diagnostic path dereferenced an empty `std::optional` (`currStmtSource()`). That
handler runs during the direct statement walk in `ResolveSpecificationParts`,
where the statement source is not set, so the dereference aborted the compiler.
The diagnostic source is now derived from the prefix, and the same fix is applied
to every diagnostic in the `LAUNCH_BOUNDS()` and `CLUSTER_DIMS()` handlers.
When only the maximum-threads-per-block operand is present, the
`cuf.launch_bounds` attribute omits `minBPM` (now an optional parameter) and no
`nvvm.minctasm` attribute is generated.
Assisted-by: AI
Add a new MLIR pass that delays cuf.alloc operations for allocatable
descriptor (box) types from function entry to just before their first
use. This defers cudaMallocManaged calls so that programs can call
cudaSetDevice or initialize communication libraries before any CUDA
context is created.
Added support for intrinsics that query the coarray in argument:
- Adding lowering and operation in MIF Dialect for UCOBOUND, LCOBOUND,
COSHAPE and IMAGE_INDEX
- Adding support of coarray argument for THIS_IMAGE in MIF Dialect (and
the lowering)
---------
Co-authored-by: Dan Bonachea <dobonachea@lbl.gov>
Co-authored-by: jeanPerier <jean.perier.polytechnique@gmail.com>
…bytes
in https://github.com/llvm/llvm-project/pull/194969 the default array
alignment was changed to 64bytes for Fortran globals. In this PR the
malloc wrapper in the flang runtime is modified to accept an alignment
argument (0 default) and for cases above 16 byte (std::max_align_t) uses
`aligned_alloc` instead of malloc.
The issue of alignment was discussed in this RFC:
https://discourse.llvm.org/t/rfc-alignment-of-global-arrays/90397/13
Assisted-by: Opus 4.6
An `!$acc routine` directive placed directly within an interface block (as an
interface-specification, e.g. preceding the interface body it names) failed to
parse and produced a misleading cascade of errors.
Accept the OpenACC ROUTINE directive as an additional interface-specification
alternative, mirroring the existing extension that allows a ROUTINE directive in
a module subprogram part / at the top level. The grammar is widened narrowly to
the ROUTINE construct only (not the whole OpenACC declarative construct). A named
directive applies to the interface body it names; name resolution already runs in
a separate pass after the interface-body symbols are created, so the
directive-before-body ordering resolves correctly and the ROUTINE information is
attached to the named procedure's symbol.
This is documented as a Flang extension in OpenACC-extensions.md.
Assisted-by: AI
This PR completes the addition of enumeration types to the symbol table.
It adds core semantic analysis for ENUMERATION TYPE, name resolution,
expression analysis for constructors and relational operators, and
SELECT CASE support. Also includes a bug fix for non-constant
enumeration comparisons in Relate() (scope() → GetScope()).
It is the 2nd of 5 stacked PRs.
AI Usage Disclosure: AI tools (Claude Opus 4.6) were used to assist with
implementation of this feature and test code generation. I have
reviewed, modified, and tested all AI-generated code.
---------
Co-authored-by: Kevin Wyatt <kwyatt@hpe.com>
The function filtering pass, intended to remove host-only functions from
a target device MLIR module on an OpenMP offloading application, also
removes any remaining host operations placed within functions containing
an `omp.target` region.
One problem with this approach is that the function filtering pass runs
relatively early, so followup passes can end up introducing new host
operations that would not be caught in advance. Specifically, there is a
known case where the `LowerWorkdistribute` pass can add
`omp.target_data` operations that then trigger an assertion during MLIR
to LLVM IR translation.
By splitting the host op deletion logic into its own pass, we can delay
it to run after `LowerWorkdistribute`, preventing these transformations
from producing an invalid target device MLIR module. The logic of the
new pass and associated test have been tweaked to properly handle FIR.
Fixes#180503
Instead of using std::get_if or parser::Unwrap, use utility functions.
This also adds two additional functions: GetLocatorFromObj and
GetCommonBlockFromObj.
Add an optional `invariant` unit attribute to `fir.load`, modeled on the
existing `nontemporal` attribute. When set, it indicates the referenced
memory location holds the same value at all points in the program where
it is dereferenceable, so the load may be treated as invariant.
The FIR-to-LLVM LoadOpConversion forwards the attribute to the invariant
flag of `llvm.load` (which lowers to `!invariant.load`).
`!invariant.load` metadata might be useful for OpenACC/OpenMP offload,
e.g. by using MLIR AA we can mark `fir.load`s inside the GPU kernels
letting the target codegen to generate read-only data cache load
instructions.
Under -gpu=mem:unified, the CUDA data attribute compatibility check
(AreCompatibleCUDADataAttrs) allows an unattributed actual argument to
bind to a Device/Managed/Unified dummy. This relaxation is correct for
variables whose storage is accessible from device code under unified
memory, but not for non-variable expression results (e.g. intrinsic call
results like RESHAPE(...)), which are host temporaries.
Without this fix, a generic SUM call like sum(reshape(hostArray, ...),
dim=2) could resolve to a device-specific overload instead of the
intrinsic, because the RESHAPE result was incorrectly considered
compatible with a DEVICE-attributed dummy.
This patch:
- Adds an actualIsVariable parameter to AreCompatibleCUDADataAttrs
(defaulting to true for backward compatibility) and gates the
unified-memory relaxation for Device/Managed/Unified dummies on it.
- Passes actualIsVariable from CheckExplicitDataArg in check-call.
- Stops setting actualCanUseImplicitCudaMemoryMode for non-variable
expressions in GetMatchingDistance, so the CUDA matching distance
ranking is also correct.
…t firstprivate
In e57ebfdaef the flag
enableDelayedPrivatizationStaging was removed, enabling
Delayed Privatization for firstprivate on OpenMP target regions,
this resulted in regressions for implicit firstprivate tests:
This commit readds enableDelayedPrivatizationStaging flag for the implicit
case and also guards the explicit change via
`enableDelayedPrivatization` instead of enableDelayedPrivatizationStaging.
Change OpenACC `DEFAULT(NONE)` scalar handling to use the
pre-OpenACC-3.2 scalar behavior by default while emitting a warning.
Scalars referenced in a `default(none)` compute region without an
explicit data clause now warn by default instead of erroring. Arrays and
other non-scalars still error under `default(none)`.
Users can opt into OpenACC 3.2 strict scalar behavior with:
`-fopenacc-default-none-scalars-strict` and the default scalar warning
can be suppressed with: `-Wno-openacc-default-none-scalars-strict`
This small code below would trigger not easy to understand error like
`error: expected 'END'`. Ignore such directive and emit the standard
warning.
```
module acc_routine_bad
contains
!$acc routine
subroutine acc_routine20()
end subroutine
end module
```
Added support for inlining hlfir.assign when both LHS and RHS are slices
of the same array. When overlap between the slices cannot be determined,
the pass introduces a disjointness check:
- genIndexBasedDisjointnessCheck(..) is used when both sides are
sections of the same array.
- genAddressBasedDisjointnessCheck(..) used as a fallback for more
complex cases.
At runtime:
- If the slices are disjoint, a direct element-wise copy is performed
without allocating a temporary buffer.
- If overlap is possible, a temporary buffer is allocated, the RHS is
first copied into it and then the data is copied from the temporary
buffer to the LHS.
Fixes https://github.com/llvm/llvm-project/issues/203228
Added support for option `-fpseudo-probe-for-profiling` in flang.
- When the option `-fpseudo-probe-for-profiling` is passed, the compiler
sets the` PseudoProbeForProfiling` flag and triggers the
`SampleProfileProbePass`. This pass inserts `llvm.pseudoprobe(..)`
intrinsic calls and `!llvm.pseudo_probe_desc` metadata into the IR.
Allow function call references and reserved locator names as parts of
OmpObject. Function calls and array element accesses have the same
syntax, and the OmpObject parser will parse them as function calls. This
is then corrected (if needed) immediately after the name resolution is
complete.
There are no clause-specific semantic checks of proper locators.
Existing code will check if a proper locator is specified on a clause
that allows it.
Lowering of proper locators to MLIR is not implemented, and a TODO
message is emitted.
Replace the lowering-only `makeVariantMatchInfo` helper with a single
shared `semantics::omp::MakeVariantMatchInfo`. It builds the
VariantMatchInfo from a parsed context selector and returns the optional
non-constant user condition (as before). Update metadirective lowering
to use it and drop the duplicated Lower/OpenMP copy.
Selector features that variant selection cannot yet honour
(target_device selectors, and clause/extension trait properties) are not
match-info concerns, so they are kept out of `MakeVariantMatchInfo`.
Detection lives in a separate, pure helper
`FindUnsupportedSelectorFeature`; the caller diagnoses the feature in
its own terms (metadirective lowering emits a TODO) before building the
match info. `MakeVariantMatchInfo` checks the precondition. NFC for
metadirective.
Co-authored-by: Cursor
---------
Co-authored-by: Cursor <cursoragent@cursor.com>
Lowering is generating patterns when forwarding OPTIONAL in calls that
looks like:
```
%present = fir.is_present %var : (T) -> i1
%if_result = fir.if %present -> (T) {
fir.result %var : T
} else {
%absent = fir.absent T
fir.result %absent : T
}
```
This specific pattern is a no-op and `%var` can be used directly. The
lowering logic that generates such patterns is inside non trivial
compiler code that has to deal with more complex scenarios where the
code inside the fir.if is more complex. Add a FIR pattern to
canonicalize such code to help with later analysis (like aliasing).
Tom Eccles
Jean-Didier PAILLEUX
Sergey Shcherbinin
Jason Van Beusekom
Abid Qadeer
jeanPerier
Matsu
Slava Zakharin
Ivan R. Ivanov
Susan Tan (ス-ザン タン)
Ferran Toda
Caroline Newcombe
Matt Arsenault
Sergio Afonso
Zhen Wang
Jack Styles
Eugene Epshteyn
Razvan Lupusoru
Chi-Chun, Chen
Valentin Clement (バレンタイン クレメン)
Ron Green [NVIDIA]
kwyatt-extandKevin Wyatt
Krzysztof Parzyszek
Andre Kuhlenschmidt
Kaviya Rajendiran
Akash Banerjee
