Fortran 2023 7.6.1 errata f23/013 changes R762 so an interoperable
enumerator may be a named-constant initialized by a
boz-literal-constant, and specifies the value as
INT(boz-literal-constant, C_INT).
Handle that named-constant restriction directly in enum resolution while
preserving the usual scalar integer expression checks for non-enumerator
contexts. This mirrors the general named-constant initializer path,
which already converts BOZ values through the declared type before
folding.
The problem is that the automatic object was used as the
threadprivate argument.
Solution:
From the documentation (J3/24-007):
> C814 An automatic data object shall not have the SAVE attribute.
> C862 The SAVE attribute shall not be specified for a dummy argument,
> a function result, an automatic data object, or an object
> that is in a common block.
The bare SAVE statement makes any valid local variables tagged
with the SAVE attribute. As the restriction mentioned above, the
automatic object cannot have a SAVE attribute and shall not be
used as the THREADPRIVATE argument. This patch adds semantic
checks for dummy arguments, function results, and automatic data
objects in both THREADPRIVATE and DECLARE TARGET directives.
[flang][OpenMP] Fix DECLARE TARGET to reject distinct result name
When a function uses a distinct result name (e.g. result(res)),
DECLARE TARGET(res) was not diagnosed because ResolveOmpDesignator
rebound it to the procedure, skipping the IsFunctionResult check.
The check added for F2023 19.4 p2 (statement function dummy argument
name may be the same as an accessible name only if that name is a scalar
variable) looked the name up through the whole host scope chain and,
failing that, in the global scope. Both lookups overshoot what is
"accessible" in the scoping unit:
* Per F2023 19.5.1.4 p2 item (11), the appearance of a name as a
dummy-arg-name in a stmt-function-stmt makes any host entity of that
name inaccessible by host association throughout the scoping unit, so a
host entity can never conflict with the dummy.
* A global entity to which the scoping unit makes no reference at all is
not accessible in it; otherwise conformance would depend on whether
unrelated program units happen to be compiled in the same file.
This caused bogus errors on conforming code such as
```
program p
logical, external :: x
contains
subroutine s
real :: f
f(x) = x * x ! 'x' is an implicitly typed real dummy here
end
end
```
where the dummy 'x' must be implicitly typed real (both NAG and gfortran
accept this and type it that way; flang already typed it correctly, but
rejected the program).
Restrict the lookup to names made visible by the scoping unit itself
(declared, referenced, or USE-associated there), matching the lookup
HandleStmtFunction already uses to type the dummies.
Assisted-by: AI
This change builds on top of the work done in PR #203959
- Diagnose interop destroy without an interop variable
- Lower array-element / derived-component interop-vars via their
designator
- Require interop-var to be a scalar integer of omp_interop_kind
- Remove "hsa" prefer_type (no omp_ifr_hsa in the runtime)
- Fix a latent bug in Parser/OpenMP/interop-construct.f90
Statement functions are inlined during lowering, so they should not be
treated as host procedure calls. Recursively validate their expressions
to reject any host procedure calls within the body.
In `HandleAttributeStmt`, when an `ASYNCHRONOUS` or `VOLATILE` statement
names a variable not already in the current scope, flang creates a
host-association symbol — but only for `Subprogram` and `BlockConstruct`
scopes. Submodule scopes have kind `Module`, so the check excluded them.
This caused a fresh `EntityDetails` symbol (zero-initialized) to be
created instead of a `HostAssocDetails` symbol pointing at the ancestor
module variable, producing value 0 instead of the initialized value.
Fix: extend the condition to also fire when `currScope().IsSubmodule()`.
`Scope::FindSymbol` already traverses into the parent module for
submodules, so the host variable is found correctly.
Note: the `volatile`/`asynchronous` `fortran_attrs` are not propagated
to `hlfir.declare` for host-associated variables in general
(subprograms,
BLOCK constructs, and submodules); that is a separate lowering issue
tracked in #208588.
Fixes#208362
Assisted-By: AI
`AccAttributeVisitor::Post(Name)` "adjusts" every name referenced inside
an OpenACC region to the symbol found in the current scope. A construct
entity -- a `DO CONCURRENT`/`FORALL` index-name, or a variable declared
in a `BLOCK` construct -- lives in its own scope nested within the
region, but the visitor's current scope does not descend into that
construct scope. `FindSymbol` therefore resolves the name to a
like-named variable in an enclosing scope and rebinds the reference to
it.
When such an entity shadows an enclosing DO variable and the loop is
associated with an OpenACC construct, this makes the entity alias that
active DO variable, so referencing or defining it wrongly triggers
"Cannot redefine DO variable" -- even though the code is legal and
compiles without error when OpenACC is disabled. This affected a `DO
CONCURRENT` index-name as well as a variable declared in a `BLOCK`
construct nested in the loop, e.g.:
```fortran
do i = 1, m
!$acc parallel loop gang vector collapse(2) independent
do concurrent (j = 1:n, ii = 1:m)
block
integer :: i ! shadows the enclosing `do i`
i = ii
a(i,j) = b(i,j)
end block
end do
end do
```
Skip the adjustment when the resolved symbol is owned by a scope
contained in the current scope, i.e. it was declared within the region;
name resolution already bound it correctly. The `DO CONCURRENT`
induction variables are then privatized as their own construct-local
entities rather than the shadowed enclosing variable, as reflected in
the updated `acc-loop.f90` lowering test.
Add a Semantics regression test covering the `DO CONCURRENT` and `BLOCK`
cases.
Flang does not enforce `DEFAULT(NONE)` for variables referenced in a
loop
associated with a standalone METADIRECTIVE. Consequently, it accepts
invalid
programs when a potentially selectable loop variant specifies
DEFAULT(NONE) but
omits a referenced variable from its data-sharing clauses. This change
checks the
associated loop for each potentially selectable `DEFAULT(NONE)` variant
and
diagnoses missing data-sharing attributes as it does for ordinary OpenMP
loop
constructs.
Standalone metadirectives and their associated loops are separate
parse-tree
nodes. For example:
```fortran
!$omp metadirective &
!$omp& when(implementation={vendor(llvm)}: &
!$omp& parallel do default(none) shared(n, a)) default(nothing)
do i = 1, n
a(i) = x
end do
```
Unlike an ordinary `PARALLEL DO`, the loop is not nested under the
directive:
```text
METADIRECTIVE
`-- WHEN
`-- PARALLEL DO DEFAULT(NONE) SHARED(n, a)
DO
`-- a(i) = x
```
Name resolution therefore leaves the variant context before visiting the
loop and
does not diagnose the missing data-sharing attribute for `x`.
After associating a metadirective with its loop, validate DEFAULT(NONE)
for each
potentially selectable loop variant. Preserve the existing rules for
predetermined and explicit data-sharing attributes, static locals,
nested
constructs, specification expressions, and statically inapplicable
variants.
Assisted with codex.
## [flang][OpenMP] Implement the `error` directive
### Summary
Implements the OpenMP `error` directive (OpenMP 5.1, 2.5.4) in Flang.
The directive prints a message and, with `severity(fatal)`, stops.
Handling depends on the `at` clause:
- **`at(compilation)`** (the default): handled in semantics. Emits a
compile-time warning or error and produces no IR.
- **`at(execution)`**: lowered to a new `omp.error` op, which becomes a
call to `__kmpc_error`.
Before this patch the directive was unimplemented and hit a `not yet
implemented: OmpErrorDirective` error during lowering.
Fixes#204240.
### Details
Changes go through the compiler pipeline.
**Flang semantics**
- `Enter(OmpErrorDirective)` emits the compile-time diagnostic for the
`at(compilation)` form, using the directive's severity and the `MESSAGE`
text when present.
- `Enter(OmpClause::Message)` checks that the `MESSAGE` expression is a
scalar, default-kind `CHARACTER`, which is required to pass it to the
runtime as a `const char*`.
**Flang lowering**
- `genErrorDirective` lowers the `at(execution)` form only.
- A constant message is stored in the `message` attribute. A
non-constant message is copied to a null-terminated buffer
(`genNullTerminatedString`) and passed as the `message_expr` operand,
because `__kmpc_error` expects a NUL-terminated string.
**New `omp.error` op (MLIR OpenMP dialect)**
- Adds a `ClauseSeverity` enum (`warning = 1`, `fatal = 2`, matching the
`__kmpc_error` argument) and `ClauseSeverityAttr`.
- `omp.error` takes a required `severity` attribute and an optional
message, given either as:
- `message` — a constant string attribute, or
- `message_expr` — a pointer to a null-terminated string.
- A verifier rejects providing both `message` and `message_expr`.
**MLIR → LLVM IR translation**
- Translates `omp.error` to `__kmpc_error(ident_t *, i32 severity, ptr
message)`.
**OMPIRBuilder**
- Adds `OpenMPIRBuilder::createError(...)`, which builds the call, maps
`fatal`→`2` and `warning`→`1`, and passes a null pointer when there is
no message.
### Tests
- **Parser** (`error-unparse.f90`): unparse and parse-tree output.
`at(compilation) severity(fatal)` cases that would stop compilation were
moved to `at(execution)`.
- **Semantics** (`Semantics/OpenMP/error.f90`): compilation
warning/error/default-severity cases, plus the new `MESSAGE` type, rank,
and kind errors.
- **Lowering** (`Lower/OpenMP/error.f90`): constant message, no message,
empty message, runtime `message_expr`, and `at(compilation)` cases that
must not produce `omp.error`. The old `Lower/OpenMP/Todo/error.f90` test
is removed.
- **MLIR** (`ops.mlir`, `invalid.mlir`): round-trip of each `omp.error`
form and the both-messages verifier error.
- **LLVM IR** (`openmp-llvm.mlir`): `__kmpc_error` calls for warning and
fatal, constant and runtime messages, and the null-message cases.
Assisted-by: Github Copilot
After https://github.com/llvm/llvm-project/pull/211269 a generic
interface name is rejected in acc routine. This PR relaxes this by
allowing it when the interface name and the specific are the same - and
attaches the information to the specific itself.
OpenACC ties the named argument of a ROUTINE directive to a subroutine
or function. A generic interface name is therefore not a valid target.
Thus adding an explicit check for this plus informational error message.
This adds support for the OpenMP interop construct, which enables
interoperability with foreign runtime environments (e.g., CUDA, HIP).
The init, use, and destroy action clauses are lowered from the Fortran
parse tree to MLIR interop ops and translated to the appropriate runtime
calls. The omp_lib module is extended with interop query API interfaces,
property constants, and return code constants. A semantics fix relaxes
the depend/targetsync constraint for use/destroy directives where the
interop type was established at a separate init site. Unsupported depend
clauses and prefer_type selector-list forms emit explicit diagnostics
instead of being silently dropped.
Assisted-by: Claude Opus 4.6
Reland of combined #208528 and #210753.
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.
Diagnose two DECLARE VARIANT restrictions from the OpenMP specification
(5.2 [7.5], 6.0 [9.6]) that were previously accepted without error:
- If a procedure is determined to be a function variant through more
than one DECLARE VARIANT directive, the construct selector set of their
context selectors must be the same.
- A procedure determined to be a function variant may not be specified
as a base function in another DECLARE VARIANT directive.
Assisted-by: Cursor
A missing `parallel_loop` directive in the `nestedTeamsAllowedSet`
currently causes code like the following to report a semantics issue:
```f90
!$omp teams
!$omp parallel loop
do i=1, 10
end do
!$omp end teams
```
This patch makes sure to accept it, since `parallel` is an allowed
construct to be immediately nested inside of `teams`, even when combined
with another construct.
Extend the loop/routine-level compatibility check to `FunctionReference`
calls. Resolve `!$acc routine` information from procedure pointers and
interface symbols, not only `SubprogramDetails`.
Lower a reduction clause naming a USE-associated, renamed, or merged-generic
user-defined reduction, which currently aborts lowering with a "not yet
implemented" TODO. Semantics already accepts these programs (since #200329).
This covers the defined-operator, intrinsic-operator, and named forms, each for
one or several types in a single declare reduction. One shape is still deferred.
A reduction on an allocatable or pointer operand of a trivial intrinsic element
type (integer, real, complex, or logical) is not lowered yet and reaches a clean
TODO, tracked by #186765. Character and derived-type allocatable reductions do
lower.
Fixes#207255.
Fixes#207489.
An imported reduction is materialized lazily. Its `omp.declare_reduction` op is
created on demand, only when a reduction clause has resolved the exact reduction
it references, through the same lowering path a local declare reduction uses.
There is no separate eager pass that walks accessible scopes to pre-create ops.
The module only-list property is therefore structural. A program using
`use m, only: <non-reduction>` never emits m's reduction, because nothing
references it. One path handles every import form: plain and renamed USE, nested
and merged generic operators, submodules, and intrinsic operators.
Resolve the use-site reduction to its source symbol through the shared semantic
resolver (`Fortran::semantics::omp::FindUserReductionSymbol`), extracted from
check-omp-structure.cpp into Semantics/openmp-utils so lowering and the semantic
checks share one resolution authority. Name the op from the resolved symbol's
ultimate name and owning scope via `getScopedUserReductionName`, on both the
clause and directive sides so the two cannot drift.
For a multi-type declare reduction, the same-file directive emits one op per
listed type and binds each variable to the op for its type. Lazy materialization
is stricter, emitting only the per-type op a clause requests, so an unsupported
sibling type does not block a supported one. The per-type lists advance in
lockstep: type, combiner, and initializer. The type suffix is unconditional,
including for a single-type reduction, so a multi-type per-type op cannot collide
with a single-type op sharing the same type-suffixed spelling.
Intrinsic-operator user reductions are named by their defining scope and ultimate
symbol, matching the defined-operator and named forms. Previously they used a
global name derived from the operator and element type, so two declare reductions
for the same operator and type in different scopes (two modules, or two
procedures in one file) could collide onto one op and silently bind the wrong
combiner. Builtin intrinsic reductions keep their global names.
For separate compilation the imported directive carries bound names but no typed
expressions, because mod-file reading runs only ResolveNames. The materializer
therefore runs expression analysis over the directive before lowering. Creating
the op from an imported module directive during lowering follows the imported
declare-mapper precedent.
An operator-less declare reduction (a special function such as max or min, or an
intrinsic operator such as `+` on an intrinsic type) has an internal mangled name
that is not valid Fortran, and no re-exported operator to recover through. When
such a reduction is re-exported through a facade, emitting it as a `use, only`
item produced a module file that failed to re-parse. The module-file writer
instead re-exports the defining module with a plain `use`, because a mangled name
cannot appear in an only-list. Entities the facade renames or makes private keep
their own use, rename, or private items, so the plain `use` does not widen their
visibility. The reduction then arrives as one shared use-association owned by the
defining module, binding the user's combiner rather than a facade-owned duplicate
that would fall back to the intrinsic. A private reduction is not re-exported.
When a reduction identifier resolves to more than one distinct user-defined
reduction for the list item's type, the clause is rejected as ambiguous rather
than selecting a combiner silently by USE order. Distinctness is by canonical
identity, the defining module and reduction name, so one reduction reached
through several USE, rename, or facade paths collapses to one candidate while
reductions from different modules stay separate. One case is not distinguished.
A same-named module loaded at two versions (a stale hermetic embed versus a
rebuilt copy) collapses to one candidate and is still resolved by USE order,
matching the prior behavior.
Assisted-by: Claude Opus 4.8, GPT-5.5.
This PR
- Adds semantic analysis for F2023 conditional arguments (F2023
R1526-R1528), enforcing constraints F2023 C1538-C1545.
- Adds shape handling and folding for conditional args.
- Adds semantic tests.
(More details in commit messages)
This PR does not implement lowering to HLFIR; That will be a separate
PR.
AI use disclaimer: The changes in this PR were substantially
generated with the assistance of AI (claude opus 4.6 via Github
CoPilot). In accordance with the LLVM project’s AI use policy, I have
reviewed and tested the code to the best of my ability.
This is a follow-up to https://github.com/llvm/llvm-project/pull/207088.
Post-merge review feedback:
* Rename mayVariantBeSelected to MayVariantBeSelected and move it to
the shared OpenMP semantic utilities so declare-variant checking can
reuse it.
* Explain why a WHEN selector is recorded only when its modifier list
contains exactly one element.
* Diagnose loop-associated variants that end a declaration-only
program unit, such as a module with no execution part.
Additional correctness fixes:
* Handle consecutive metadirectives and modules or submodules with
contained procedures without silently dropping an earlier
loop-associated variant.
* Prevent a loop-associated variant in an interface, BLOCK construct,
conditional branch, or nested body from being checked against an
unrelated loop outside that scope.
* Avoid false or missing loop diagnostics when match_any or match_none
combines target selectors with user or construct selectors.
Assisted with codex.
OpenMP 6.0, 9.8 "declare_simd Directive", restricts the placement of the
directive:
Any declare_simd directive must appear in the specification part of a
subroutine subprogram, function subprogram, or interface body to which
it applies.
Flang did not enforce this, and accepted the directive in the
specification part of a module, submodule, main program or block data.
Note: the restriction requiring the argument to name the procedure to
which the directive applies is not implemented here.
Fixes#205478
These tests pin currently-correct behaviors where the KIND type
parameter of a result must be preserved by constant folding, semantics,
and lowering. They all pass at head; their purpose is to guard against
silent kind loss during future refactoring of the Evaluate library's
type-parameterized machinery (see the discussion on PR #206907, where an
experimental rewrite passed check-flang while dropping kinds in several
places).
Folded LOGICAL operations: `.NOT./.AND./.OR./.EQV./.NEQV.` keep their
operand kind; ANY/ALL/PARITY and logical DOT_PRODUCT results have the
MASK/argument kind (F2023 16.9.14/16.9.16/16.9.148), including for empty
masks where there is no element value to take a kind from; MERGE and
TRANSFER keep the TSOURCE/MOLD kind; STORAGE_SIZE of folded
non-default-kind logicals folds to the right value; the default BOUNDARY
fill element of a folded EOSHIFT has the array's kind; relational
results are default logical regardless of operand kinds; and generic
resolution over LOGICAL kinds is unchanged by folding
(flang/test/Evaluate/fold-kind-logical.f90 and an extension of
flang/test/Evaluate/logical-args.f90).
Empty typed array constructors and zero-size PACK/SPREAD/RESHAPE/UNPACK
results keep the kind (and character length) of their type-spec (F2023
7.8), including across a module file into a separately compiled
consumer, and empty non-default-kind character array constructors lower
with the correct `!fir.char` element type
(flang/test/Evaluate/fold-kind-zero-size.f90,
flang/test/Semantics/modfile84.F90,
flang/test/Lower/HLFIR/array-ctor-character-kind.f90).
A substring of a character literal has the kind of the literal (F2023
9.4.1): KIND() and LEN() fold correctly, and concatenation with a
same-kind literal is accepted
(flang/test/Evaluate/fold-kind-substring.f90,
flang/test/Semantics/substring-literal-kind.f90).
Assisted-by: AI
Stop writing a module file that cannot be re-parsed (reading it crashes) when a
module re-exports a USE-associated user-defined operator declare reduction.
Such a reduction is represented by an internal symbol whose name is the mangled
operator (e.g. "op.remote." or "op.+"), which is not valid Fortran. The
module-file writer emitted it as a `use MODULE, only: op.remote.` item, which the
reader cannot parse. This affected both a plain re-export facade and an embedded
module in a hermetic module file.
Skip the use-only item for a reduction whose operator is itself re-exported (a
defined operator, or an intrinsic operator with a user interface): the reduction
is re-exported implicitly with that operator, which is emitted, and the reduction
resolver recovers it from there (FindUserReductionSymbol). This covers a
derived-type `operator(+)` reduction as well as a defined operator like
`.remote.`.
The test checks the re-exporting module file round-trips and that a consumer
reading it resolves the reduction through the operator.
Assisted-by: Claude Opus 4.8, GPT-5.5.
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
Fixes#191343
Replaced the error with warning. Now it permits calling an external
procedure with implicit interface and passing a volatile argument. There
will be a warning that the procedure should have an explicit interface.
In addition to fixing the VOLATILE error, I updated the semantic check
for the ASYNCHRONOUS attribute. Both attributes were incorrectly
throwing a hard error when used as actual arguments in an implicit
interface call. According to the Fortran 2018 standard, an explicit
interface is only mandatory when the dummy argument has the VOLATILE or
ASYNCHRONOUS attribute. Since the standard doesn't restrict actual
arguments in this scenario, I downgraded both to emit a
-Wimplicit-interface-actual warning instead.
A loop-associated metadirective variant (`do`, `simd`, ...) is only
resolved during lowering, so it is never checked as a loop construct
during semantic analysis. A malformed or non-canonical associated nest
therefore reaches lowering, which assumes a canonical nest.
This patch validates the nest that follows such a variant (the next
executable construct) during semantics, reusing the diagnostics of a
real loop-associated construct. Each applicable variant is checked
against it:
* Canonical loop: the affected loop must be a canonical DO loop, so a
`DO WHILE`, a pre-6.0 `DO CONCURRENT`, or a `DO` without loop control is
rejected.
* Nest depth: `collapse(n)` and `ordered(n)` must not exceed the depth
of the associated loop nest.
* Rectangularity: loops that must be rectangular (e.g. under `tile`) may
not have bounds that depend on an outer loop's variable.
A loop-associated variant may also have no loop nest at all. This
happens when the metadirective ends the execution part or a non-loop
construct follows it. Such a variant is rejected like a bare loop
directive with no associated loop.
Only variants that could be selected are checked: when a target is
known, a WHEN variant whose device or implementation selector cannot
match it is skipped (runtime `user` conditions never cause a skip).
Part of the feature work for #188820. Builds on top of
[#193664](https://github.com/llvm/llvm-project/issues/193664) and
[#194424](https://github.com/llvm/llvm-project/pull/194424).
Assisted with Copilot.
Fixes [#205516 ](https://github.com/llvm/llvm-project/issues/205516)
#### Problem
Per OpenMP 4.5 Sec 2.13.8 Restrictions:
> An ordered region arising from an ordered construct with both the simd
and
> threads clauses must be closely nested inside a loop SIMD region.
Flang was missing this check. `!$OMP ORDERED THREADS SIMD` inside a
standalone `!$OMP SIMD` region was accepted without diagnostic.
#### Fix
`check-omp-structure.cpp`: Added a new check in
`ChecksOnOrderedAsBlock()`. If an ORDERED directive has both SIMD and
THREADS clauses and is nested inside a standalone SIMD region (but not a
DO SIMD region), emit the following error:
`An ORDERED directive with SIMD and THREADS clauses must be closely
nested in a worksharing-loop SIMD region.`
`ordered02.f90`: Added two test cases:
Negative test: ORDERED THREADS SIMD inside !$OMP SIMD → expected error.
Positive test: ORDERED THREADS SIMD inside !$OMP DO SIMD ordered → no
error, as allowed by OpenMP 2.13.8.
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.
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
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
Implement F2023 19.4.2 conformance check for statement function dummy
argument names
Fixes: #203820
F2023 19.4.2 permits a statement function dummy argument to share its
name with an **accessible** identifier only when that identifier is a scalar
variable. Flang previously accepted all such clashes silently.
```fortran
subroutine s
real :: arr(10)
real :: ok
f(arr) = arr + 1 ! error: arr is an array
h(ok) = ok + 1 ! ok: scalar variable shadowing is permitted
end subroutine
```
Assisted-by: Claude
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
ResolveAccObject skipped DSA registration when
GetDesignatorNameIfDataRef returned nullptr (i.e. for subscripted
designators like copyin(a(1:n))). The DEFAULT(NONE) post-visitor then
found the base variable unregistered and emitted a false error even
though it was explicitly listed in a data clause.
Unify the path using GetFirstName, which extracts the base symbol from
both bare data references and array sections. Continue not registering component references. This will need to be followed up with more accurate designator checking in the near future.
An ATTRIBUTES(HOST,DEVICE) subprogram is compiled for the host as well
as the device, so referencing a host procedure (typically guarded at run
time by a test such as ON_DEVICE()) is legitimate in its host
compilation. Two semantic checks were wrongly rejecting such code:
* CheckSubprogram flagged an interface body declared inside a device
subprogram as "may not be an internal procedure of CUDA device
subprogram". An interface body (like a dummy procedure) is not an
internal procedure, so exclude both from this check. A real internal
procedure is still rejected.
* DeviceExprChecker rejected calls to host procedures with "may not be
called in device code" even from host,device subprograms. Thread the
existing isHostDevice flag of DeviceContextChecker down to
DeviceExprChecker and suppress only that diagnostic when the enclosing
subprogram is host,device. Pure device and global subprograms continue
to reject host calls, and all other device-code diagnostics are
preserved.
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>
1. Require the base and variant procedures to have compatible
characteristics as required by OpenMP 6.0 section 9.6.
2. Reject adjust_args/append_args as not yet implemented.
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`
