Tobias Burnus [Thu, 5 Jun 2025 08:36:21 +0000 (10:36 +0200)]
gcn: Update --with-arch= for newer archs
Replace hard-coded list of supported devices by directly checking
config/gcn/gcn-devices.def.
gcc/ChangeLog:
* config.gcc (--with-{arch,tune}): Use .def file to validate gcn
processor names.
* doc/install.texi (amdgcn*-*-*): Update list of devices supported
by --with-arch/--with-tune.
Patrick Palka [Fri, 6 Jun 2025 13:34:17 +0000 (09:34 -0400)]
libstdc++: Fix flat_map::operator[] for const lvalue keys [PR120432]
The const lvalue operator[] overload wasn't properly forwarding the key
type to the generic overload, causing a hard error for const keys.
Rather than correcting the forwarded type this patch just makes the
non-template overloads call try_emplace directly instead. That way we
can remove the non-standard same_as constraint on the generic overload
and match the spec more closely.
PR libstdc++/120432
libstdc++-v3/ChangeLog:
* include/std/flat_map (flat_map::operator[]): Make the
non-template overloads call try_emplace directly. Remove
non-standard same_as constraint on the template overload.
* testsuite/23_containers/flat_map/1.cc (test08): New test.
Reviewed-by: Tomasz Kamiński <tkaminsk@redhat.com> Reviewed-by: Jonathan Wakely <jwakely@redhat.com>
(cherry picked from commit 91ed3248ce26aaaee4d7471aa4edbc07b3f1a90e)
Tomasz Kamiński [Wed, 4 Jun 2025 09:05:11 +0000 (11:05 +0200)]
libstdc++: Fix format call in formatting with empty specs for durations.
This patches fixes an obvious error, where the output iterator argument was
missing for call to format_to, when duration with custom representation types
are used.
libstdc++-v3/ChangeLog:
* include/bits/chrono_io.h (__formatter_chrono:_M_s): Add missing
__out argument to format_to call.
Reviewed-by: Jonathan Wakely <jwakely@redhat.com> Signed-off-by: Tomasz Kamiński <tkaminsk@redhat.com>
(cherry picked from commit ac0a04b7a254fb8e1d8d7088336bcb4375807b1e)
Richard Biener [Fri, 30 May 2025 12:11:47 +0000 (14:11 +0200)]
tree-optimization/120357 - ICE with early break vectorization
When doing early break vectorization of a loop with a conditional
reduction the epilog creation code is confused as to before which exit
to insert the conditional reduction induction IV update. The
following make sure this is done before the main IV exit.
PR tree-optimization/120357
* tree-vect-loop.cc (vect_create_epilog_for_reduction): Create
the conditional reduction induction IV increment before the
main IV exit.
* gcc.dg/vect/vect-early-break_136-pr120357.c: New testcase.
Richard Biener [Fri, 9 May 2025 06:38:45 +0000 (08:38 +0200)]
rtl-optimization/120182 - wrong-code with RTL DSE and constant addresses
RTL DSE forms store groups from unique invariant bases but that is
confused when presented with constant addresses where it assigns
one store group per unique address. That causes it to not consider
0x101:QI to alias 0x100:SI. Constant accesses can really alias
to every object, in practice they appear for I/O and for access
to objects fixed via linker scripts for example. So simply avoid
registering a store group for them.
PR rtl-optimization/120182
* dse.cc (canon_address): Constant addresses have no
separate store group.
Richard Biener [Wed, 30 Apr 2025 09:52:17 +0000 (11:52 +0200)]
tree-optimization/120003 - missed jump threading
The following allows the entry and exit block of a jump thread path
to be equal, which can easily happen when there isn't a forwarder
on the interesting edge for an FSM thread conditional. We just
don't want to enlarge the path from such a block.
PR tree-optimization/120003
* tree-ssa-threadbackward.cc (back_threader::find_paths_to_names):
Allow block re-use but do not enlarge the path beyond such a
re-use.
* gcc.dg/tree-ssa/ssa-thread-23.c: New testcase.
* gcc.dg/tree-ssa/ssa-dom-thread-7.c: Adjust.
The following addresses a too conservative sanity check of SLP nodes
we want to promote external. The issue lies in code generation
for such external which relies on get_later_stmt to figure an
insert location. But get_later_stmt relies on the ability to
totally order stmts, specifically implementation-wise that they
are all from the same BB, which is what is verified at the moment.
The patch changes this to require stmts to be orderable by
dominance queries. For simplicity and seemingly enough for the
testcase in PR119960, this handles the case of two distinct BBs.
PR tree-optimization/119960
* tree-vect-slp.cc (vect_slp_can_convert_to_external):
Handle cases where defs from multiple BBs are ordered
by their dominance relation.
Richard Biener [Thu, 8 May 2025 08:56:16 +0000 (10:56 +0200)]
tree-optimization/116352 - amend previous fix
The previous fix restricted external vector builds to defs from
the same basic-block. That turns out too restrictive so we have
to mitigate the original issue in a different way which is
restricting it to the original case where all defs are in the
same basic-block.
PR tree-optimization/116352
* tree-vect-slp.cc (vect_build_slp_tree_2): When compressing
operands from a two-operator node make sure the resulting
operation does not mix defs from different basic-blocks.
Richard Biener [Tue, 29 Apr 2025 13:08:52 +0000 (15:08 +0200)]
tree-optimization/119960 - add validity checking to SLP scheduling
The following adds checks that when we search for a vector stmt
insert location we arrive at one where all required operand defs
are dominating the insert location. At the moment any such
failure only blows up during SSA verification.
There's the long-standing issue that we do not verify there
exists a valid schedule of the SLP graph from BB vectorization
into the existing CFG. We do not have the ability to insert
vector stmts on the dominance frontier "end", nor to insert
LC PHIs that would be eventually required.
This should be done all differently, computing the schedule
during analysis and failing if we can't schedule.
Richard Biener [Wed, 14 May 2025 14:36:29 +0000 (16:36 +0200)]
Fix regression from x86 multi-epilogue tuning
With the avx512_two_epilogues tuning enabled for zen4 and zen5
the gcc.target/i386/vect-epilogues-5.c testcase below regresses
and ends up using AVX2 sized vectors for the masked epilogue
rather than AVX512 sized vectors. The following patch rectifies
this and adds coverage for the intended behavior.
* config/i386/i386.cc (ix86_vector_costs::finish_cost):
Do not suggest a first epilogue mode for AVX512 sized
main loops with X86_TUNE_AVX512_TWO_EPILOGUES as that
interferes with using a masked epilogue.
Javier Miranda [Thu, 6 Feb 2025 09:40:57 +0000 (09:40 +0000)]
ada: Spurious compilation error with repeated loop index
When multiple for-loop statements in the same scope use the
same index name to iterate through container elements, the
compiler reports a spurious error indicating a conflict
between index names.
gcc/ada/ChangeLog:
* exp_ch7.adb (Process_Object_Declaration): Avoid generating
duplicate names for master nodes.
If the body of a loop includes a raise statement then the loop should not be
considered to be free of side-effects and therefore eligible for elimination
by the compiler.
gcc/ada/ChangeLog:
* sem_util.adb
(Side_Effect_Free_Statements): Return False if the statement list
includes an explicit (i.e. Comes_From_Source) raise statement.
The generation of the check mandated by Ada issue AI05-0073 was not done
handled properly for protected types when used through subtypes. This
patch fixes the issue.
gcc/ada/ChangeLog:
* exp_ch4.adb (Tagged_Membership): Fix for protected types.
Eric Botcazou [Fri, 24 Jan 2025 09:26:13 +0000 (10:26 +0100)]
ada: Implement built-in-place expansion of two-pass array aggregates
These are array aggregates containing only component associations that are
iterated with iterator specifications, as per RM 4.3.3(20.2/5-20.4/5).
It is implemented for the array aggregates that are used to initialize an
object, as specified by RM 7.6(17.2/3-17.3/3) for immutably limited types
and types that need finalization, but for all types like other aggregates.
gcc/ada/ChangeLog:
* exp_aggr.adb (Build_Two_Pass_Aggr_Code): New function containing
most of the code initially present in Two_Pass_Aggregate_Expansion.
(Two_Pass_Aggregate_Expansion): Remove redundant N parameter.
Implement built-in-place expansion for (static) object declarations
and allocators, using Build_Two_Pass_Aggr_Code for the main work.
(Expand_Array_Aggregate): Adjust Two_Pass_Aggregate_Expansion call.
Replace Etype (N) by Typ in a couple of places.
* exp_ch3.adb (Expand_Freeze_Array_Type): Remove special case for
two-pass array aggregates.
(Expand_N_Object_Declaration): Do not adjust the object when it is
initialized by a two-pass array aggregate.
* exp_ch4.adb (Expand_Allocator_Expression): Apply the processing
used for container aggregates to two-pass array aggregates.
* exp_ch6.adb (Validate_Subprogram_Calls): Skip calls present in
initialization expressions of N_Object_Declaration nodes that have
No_Initialization set.
* sem_ch3.adb (Analyze_Object_Declaration): Detect the cases of an
array originally initialized by an aggregate consistently.
Viljar Indus [Mon, 20 Jan 2025 13:10:22 +0000 (15:10 +0200)]
ada: Reject Valid_Value arguments originating from Standard
The constraint for Valid_Value not applying to types from Standard
should also apply to all types derived from those types.
gcc/ada/ChangeLog:
* doc/gnat_rm/implementation_defined_attributes.rst: Update the
documentation for Valid_Value.
* sem_attr.adb (Analyze_Attribute): Reject types where
the root type originates from Standard.
* gnat_rm.texi: Regenerate.
Gary Dismukes [Sat, 18 Jan 2025 01:11:12 +0000 (01:11 +0000)]
ada: Error about assignment to limited target on aggregate with "for of" iterator
The compiler reports a spurious error about an assignment to a limited
object on an aggregate of a array type with limited components that has
an association with a "for of" iterator. This is fixed by arranging to
have the Assignment_OK flag set on the indexed_names generated by the
expander for initializing the aggregate object.
gcc/ada/ChangeLog:
* exp_aggr.adb (Two_Pass_Aggregate_Expansion): Change call to Make_Assignment
for the indexed aggregate object to call Change_Make_OK_Assignment instead.
Eric Botcazou [Wed, 15 Jan 2025 19:37:48 +0000 (20:37 +0100)]
ada: Fix buffer overflow for function call returning discriminated limited record
This occurs when the discriminated limited record type is declared with
default values for its discriminants, is not controlled, and the context
of the call is anonymous, i.e. the result of the call is not assigned
to an object. In this case, a temporary is created to hold the result
of the call, with the default values of the discriminants, but the result
may have different values for the discriminants and, in particular, may
be larger than the temporary, which leads to a buffer overflow.
This problem does not occur when the context is an object declaration, so
the fix just makes sure that the expansion in an anonymous context always
uses the model of an object declaration. It requires a minor tweak to the
helper function Entity_Of of the Sem_Util package.
gcc/ada/ChangeLog:
* exp_ch6.adb (Expand_Actuals): Remove obsolete comment.
(Make_Build_In_Place_Call_In_Anonymous_Context): Always use a proper
object declaration initialized with the function call in the cases
where a temporary is needed, with Assignment_OK set on it.
* sem_util.adb (Entity_Of): Deal with rewritten function call first.
Steve Baird [Mon, 13 Jan 2025 22:18:26 +0000 (14:18 -0800)]
ada: Fix compile-time failure due to duplicated attribute subprograms.
For a given type, and for certain attributes (the 4 streaming attributes
and, for Ada2022, the Put_Image attribute), the compiler needs to keep track
of whether a subprogram has already been generated for the given
type/attribute pair. In some cases this was being done incorrectly;
the compiler ended up generating duplicate subprograms (with the same
name), resulting in compilation failures. This could occur if the prefix
of an attribute reference denoted a subtype (more precisely, a non-first
subtype). This includes the case of a subtype declaration that is implicitly
introduced by the compiler to capture the binding between a formal type
in a generic and the corresponding actual type in an instantiation.
gcc/ada/ChangeLog:
* exp_attr.adb (Expand_N_Attribute_Reference): When accessing the
maps declared in package Cached_Attribute_Ops, the key value
passed to Get or to Set should never be the entity node for a
subtype. Use the entity of the corresponding type declaration
instead.
Steve Baird [Fri, 10 Jan 2025 21:15:18 +0000 (13:15 -0800)]
ada: Avoid calling Resolve with Stand.Any_Fixed as the expected type
When we call Resolve for an expression, we pass in the expected type
for that expression. In the absence of semantic errors, that expected type
should never be any of the "Any_xxx" types declared in stand.ads (e.g.,
Any_Array, Any_Numeric, Any_Real). In particular, it should never be Any_Fixed.
Fix a case in which this rule was being violated.
gcc/ada/ChangeLog:
* sem_res.adb
(Set_Mixed_Mode_Operand): If we are about to call Resolve
passing in Any_Fixed as the expected type, then instead pass in
the fixed point type of the other operand (i.e., B_Typ).
Gary Dismukes [Fri, 10 Jan 2025 22:39:52 +0000 (22:39 +0000)]
ada: Compiler crash on array aggregate association iterating over function result
The compiler triggers a bug box when compiling an array aggregate with
an iterated_component_association that iterates over another array object,
failing when trying to retrieve a Choices field, which isn't an allowed
field for N_Iterated_Component_Association nodes. This occurs in procedure
Check_Function_Writable_Actuals, which wasn't accounting for the iterated
association forms.
gcc/ada/ChangeLog:
* sem_util.adb (Check_Function_Writable_Actuals): Add handling for
N_Iterated_Component_Association and N_Iterated_Element_Association.
Fix a typo in an RM reference (6.4.1(20/3) => 6.4.1(6.20/3)).
(Collect_Expression_Ids): New procedure factoring code for collecting
identifiers from expressions of aggregate associations.
(Handle_Association_Choices): New procedure factoring code for handling
id collection for expressions of aggregate associations with multiple
choices. Removed redundant test of Box_Present from original code.
What happens is that symtab_remove_unreachable_nodes leaves the last symbol
in kind of a limbo state: in .remove_symbols, we have:
opt7_pkg__enum_name_table/13 (Opt7_Pkg.Enum_Name_Table)
Type: variable
Body removed by symtab_remove_unreachable_nodes
Visibility: externally_visible semantic_interposition external public
References:
Referring: opt7_pkg__image/2 (read) opt7_pkg__image/2 (read)
Availability: not_available
Varpool flags: initialized read-only const-value-known
This means that the "body" (DECL_INITIAL) of the symbol has been disregarded
during reachability analysis, causing the first two symbols to be discarded:
but the DECL_INITIAL is explicitly preserved for later constant folding,
which makes it possible to retrofit the DECLs corresponding to the first
two symbols in the GIMPLE IR and ultimately leads to the crash.
gcc/
* tree-vect-data-refs.cc (vect_can_force_dr_alignment_p): Return
false if the variable has no symtab node.
gcc/testsuite/
* gnat.dg/specs/opt7.ads: New test.
* gnat.dg/specs/opt7_pkg.ads: New helper.
* gnat.dg/specs/opt7_pkg.adb: Likewise.
Harald Anlauf [Fri, 30 May 2025 17:25:15 +0000 (19:25 +0200)]
Fortran: parameter inquiries of constant complex arrays [PR102599,PR114022]
PR fortran/102599
PR fortran/114022
gcc/fortran/ChangeLog:
* expr.cc (simplify_complex_array_inquiry_ref): Helper function for
simplification of inquiry references (%re/%im) of constant complex
arrays.
(find_inquiry_ref): Use it for handling %re/%im inquiry references
of complex arrays.
(scalarize_intrinsic_call): Fix frontend memleak.
* primary.cc (gfc_match_varspec): When the reference is NULL, the
previous simplification has succeeded in evaluating inquiry
references also of arrays.
Harald Anlauf [Tue, 27 May 2025 17:23:16 +0000 (19:23 +0200)]
Fortran: fix parsing of type parameter inquiries of substrings [PR101735]
Handling of type parameter inquiries of substrings failed to due either
parsing issues or not following or handling reference chains properly.
PR fortran/101735
gcc/fortran/ChangeLog:
* expr.cc (find_inquiry_ref): If an inquiry reference applies to
a substring, use that, and calculate substring length if needed.
* primary.cc (extend_ref): Also handle attaching to end of
reference chain for appending.
(gfc_match_varspec): Discrimate between arrays of character and
substrings of them. If a substring is taken from a character
component of a derived type, get the proper typespec so that
inquiry references work correctly.
(gfc_match_rvalue): Handle corner case where we hit a seemingly
dangling '%' and missed an inquiry reference. Try another match.
Jakub Jelinek [Sat, 10 May 2025 19:20:09 +0000 (21:20 +0200)]
fortran: Fix debug info for unsigned(kind=1) and unsigned(kind=4) [PR120193]
As the following testcase shows, debug info for unsigned(kind=1)
and unsigned(kind=4) vars is wrong while unsigned(kind=2), unsigned(kind=8)
and unsigned(kind=16) look right.
Instead of objects having unsigned(kind=1) type they have character(kind=1)
and instead of unsigned(kind=4) they have character(kind=4).
This means in gdb e.g. unsigned(kind=1) :: a(2) variable initialized to
97 will print as 'aa' rather than (97, 97) etc.
While there can be just one unsigned_char_type_node and one
unsigned_type_node type, each can have arbitrary number of variants
(e.g. consider C
typedef unsigned char uc;
where uc is a variant type to unsigned char) or even distinct types
with different TYPE_MAIN_VARIANT.
The following patch uses a variant of the character(kind=4) type
for unsigned(kind=4) and a distinct type based on character(kind=1)
type for unsigned(kind=1). The reason for the latter is that
unsigned_char_type_node has TYPE_STRING_FLAG set on it, so it has
DW_AT_encoding DW_ATE_unsigned_char rather than DW_ATE_unsigned and
so the debugger then likes to print it as characters rather than numbers.
That is IMHO in Fortran desirable for character(kind=1) but not for
unsigned(kind=1). I've made sure TYPE_CANONICAL of the unsigned(kind=1)
type is still character(kind=1), so they are considered compatible by
the middle-end also e.g. for aliasing etc.
2025-05-10 Jakub Jelinek <jakub@redhat.com>
PR fortran/120193
* trans-types.cc (gfc_init_types): For flag_unsigned use
build_distinct_type_copy or build_variant_type_copy from
gfc_character_types[index_char] if index_char > -1 instead of
gfc_character_types[index_char] or
gfc_build_unsigned_type (&gfc_unsigned_kinds[index]).
Jerry DeLisle [Sat, 31 May 2025 15:57:22 +0000 (08:57 -0700)]
Fortran: Fix handling of parsed format strings.
Previously parsed strings with errors were being cached such
that subsequent use of the format string were not being
checked for errors.
PR libfortran/119856
libgfortran/ChangeLog:
* io/format.c (parse_format_list): Set the fmt->error
message for missing comma.
(parse_format): Do not cache the parsed format string
if a previous error ocurred.
Jerry DeLisle [Wed, 28 May 2025 14:56:12 +0000 (07:56 -0700)]
Fortran: Adjust handling of optional comma in FORMAT.
This change adjusts the error messages for optional commas
in format strings to give a warning at compile time unless
-std=legacy is used. This is more consistant with the
runtime library. A missing comma separator should not be
encouraged as it is non-standard fortran.
PR fortran/119856
gcc/fortran/ChangeLog:
* io.cc: Set missing comma error checks to STD_STD_LEGACY.
Patrick Palka [Thu, 29 May 2025 14:12:23 +0000 (10:12 -0400)]
libstdc++: Compare keys and values separately in flat_map::operator==
Instead of effectively doing a zipped comparison of the keys and values,
compare them separately to leverage the underlying containers' optimized
equality implementations.
libstdc++-v3/ChangeLog:
* include/std/flat_map (_Flat_map_impl::operator==): Compare
keys and values separately.
Patrick Palka [Thu, 29 May 2025 14:11:57 +0000 (10:11 -0400)]
libstdc++: Fix tuple/pair confusion with std::erase_if(flat_map) [PR120465]
std::erase_if for flat_map/multimap is implemented via ranges::erase_if
over a zip_view of the keys and values, the value_type of which is a
tuple, but the given predicate needs to be called with a pair (flat_map's
value_type). So use a projection to convert the tuple into a suitable
pair.
PR libstdc++/120465
libstdc++-v3/ChangeLog:
* include/std/flat_map (_Flat_map_impl::_M_erase_if): Use a
projection with ranges::remove_if to pass a pair instead of
a tuple to the predicate.
* testsuite/23_containers/flat_map/1.cc (test07): Strengthen
to expect the argument passed to the predicate is a pair.
* testsuite/23_containers/flat_multimap/1.cc (test07): Likewise.
Co-authored-by: Jonathan Wakely <jwakely@redhat.com> Reviewed-by: Tomasz Kamiński <tkaminsk@redhat.com> Reviewed-by: Jonathan Wakely <jwakely@redhat.com>
(cherry picked from commit 589b27ec5769410e036df57645ff1eb7c765f692)
Jerry DeLisle [Tue, 20 May 2025 02:41:16 +0000 (19:41 -0700)]
Fortran: Fix c_associated argument checks.
PR fortran/120049
gcc/fortran/ChangeLog:
* check.cc (gfc_check_c_associated): Use new helper functions.
Only call check_c_ptr_1 if optional c_ptr_2 tests succeed.
(check_c_ptr_1): Handle only c_ptr_1 checks.
(check_c_ptr_2): Expand checks for c_ptr_2 and handle cases
where there is no derived pointer in the gfc_expr and check
the inmod_sym_id only if it exists. Rephrase error message.
* misc.cc (gfc_typename): Handle the case for BT_VOID rather
than throw an internal error.
gcc/testsuite/ChangeLog:
* gfortran.dg/pr120049_a.f90: Update test directives.
* gfortran.dg/pr120049_b.f90: Update test directives
* gfortran.dg/pr120049_2.f90: New test.
* gfortran.dg/c_f_pointer_tests_6.f90: Adjust dg-error
directive.
Thomas Koenig [Fri, 30 May 2025 11:31:58 +0000 (13:31 +0200)]
Type mismatch for passed external function
This obvious and simple patch fixes a 15/16 regression where the
typespec of a global function was in the RESULT clause and not
in the symbol itself.
gcc/fortran/ChangeLog:
PR fortran/120355
* interface.cc (compare_parameter): If the global function has a
result clause, take typespec from there for the comparison against
the dummy argument.
gcc/testsuite/ChangeLog:
PR fortran/120355
* gfortran.dg/interface_62.f90: New test.
Sandra Loosemore [Mon, 26 May 2025 19:21:48 +0000 (19:21 +0000)]
OpenMP: Fix ICE and other issues in C/C++ metadirective error recovery.
The new testcase included in this patch used to ICE in gcc after
diagnosing the first error, and in g++ it only diagnosed the error in
the first metadirective, ignoring the second one. The solution is to
make error recovery in the C front end more like that in the C++ front
end, and remove the code in both front ends that previously tried to
skip all the way over the following statement (instead of just to the
end of the metadirective pragma) after an error.
gcc/c/ChangeLog
* c-parser.cc (c_parser_skip_to_closing_brace): New, copied from
the equivalent function in the C++ front end.
(c_parser_skip_to_end_of_block_or_statement): Pass false to
the error flag.
(c_parser_omp_context_selector): Immediately return error_mark_node
after giving an error that the integer trait property is invalid,
similarly to C++ front end.
(c_parser_omp_context_selector_specification): Likewise handle
error return from c_parser_omp_context_selector similarly to C++.
(c_parser_omp_metadirective): Do not call
c_parser_skip_to_end_of_block_or_statement after an error.
gcc/cp/ChangeLog
* parser.cc (cp_parser_omp_metadirective): Do not call
cp_parser_skip_to_end_of_block_or_statement after an error.
gcc/testsuite/ChangeLog
* c-c++-common/gomp/declare-variant-2.c: Adjust patterns now that
C and C++ now behave similarly.
* c-c++-common/gomp/metadirective-error-recovery.c: New.
Tobias Burnus [Wed, 28 May 2025 13:14:14 +0000 (15:14 +0200)]
libgomp.fortran/metadirective-1.f90: Expect 'error:' for nvptx compile [PR118694]
This should have been part of commit r16-838-gb3d07ec7ac2ccd or r16-883-g5d6ed6d604ff94 - all showing the same issue:
'!$omp target' followed by a metadirective with 'teams'; if
the metadirective cannot be early resolved, a diagnostic
error is shown about using directives between 'target' and
'teams'.
While the message is misleading, the problem is that the
host invokes 'target' differently when 'teams' is present;
in this case, host fallback + amdgcn offload require the
no-teams case, nvptx offload the teams case such that it
only can be resolved at runtime.
Mark the error as 'dg-bogus + xfail' to silence the FAIL,
when nvptx offloading is compiled for. (If not, the
metadirective can be resolved early during compilation.)
libgomp/ChangeLog:
PR middle-end/118694
* testsuite/libgomp.fortran/metadirective-1.f90: xfail when
compiling (also) for nvptx offloading as an error is then expected.
Eric Botcazou [Tue, 27 May 2025 17:42:17 +0000 (19:42 +0200)]
Fix IPA-SRA issue with reverse SSO on specific pattern
IPA-SRA generally works fine in the presence of reverse Scalar_Storage_Order
by propagating the relevant flag onto the newly generated MEM_REFs. However
we have been recently faced with a specific Ada pattern that it does not
handle correctly: the 'Valid attribute applied to a floating-point component
of an aggregate type with reverse Scalar_Storage_Order.
The attribute is implemented by a call to a specific routine of the runtime
that expects a pointer to the object so, in the case of a component with
reverse SSO, the compiler first loads it from the aggregate to get back the
native storage order, but it does the load using an array of bytes instead
of the floating-point type to prevent the FPU from fiddling with the value,
which yields in the .original dump file:
Of course that's a bit convoluted, but it does not seem that another method
would be simpler or even work, and using VIEW_CONVERT_EXPR to toggle the SSO
is supposed to be supported in any case (unlike aliasing or type punning).
The attached patch makes it work. While the call to storage_order_barrier_p
from IPA-SRA is quite natural (the regular SRA has it too), the tweak to the
predicate itself is needed to handle the scalar->aggregate conversion, which
is admittedly awkward but again without clear alternative.
gcc/
* ipa-sra.cc (scan_expr_access): Also disqualify storage order
barriers from splitting.
* tree.h (storage_order_barrier_p): Also return false if the
operand of the VIEW_CONVERT_EXPR has reverse storage order.
gcc/testsuite/
* gnat.dg/sso19.adb: New test.
* gnat.dg/sso19_pkg.ads, gnat.dg/sso19_pkg.adb: New helper.
Jonathan Wakely [Wed, 21 May 2025 14:29:02 +0000 (15:29 +0100)]
libstdc++: Fix vector(from_range_t, R&&) for exceptions [PR120367]
Because this constructor delegates to vector(a) the object has been
fully constructed and the destructor will run if an exception happens.
That means we need to set _M_finish == _M_start so that the destructor
doesn't try to destroy any elements.
libstdc++-v3/ChangeLog:
PR libstdc++/120367
* include/bits/stl_vector.h (_M_range_initialize): Initialize
_M_impl._M_finish.
* testsuite/23_containers/vector/cons/from_range.cc: Check with
a type that throws on construction.
exceptions during construction.
Tobias Burnus [Mon, 26 May 2025 17:50:40 +0000 (19:50 +0200)]
c-c++-common/gomp/{attrs-,}metadirective-3.c: Fix expected result [PR118694]
With compilation for nvptx enabled, two issues showed up:
(a) "error: 'target' construct with nested 'teams' construct contains
directives outside of the 'teams' construct"
See PR comment 9 why this is difficult to fix.
Solution: Add dg-bogus and accept/expect the error for 'target offload_nvptx'.
(b) The assumptions about the dump for 'target offload_nvptx' were wrong
as the metadirective was already expanded to a OMP_NEXT_VARIANT
construct such that no 'omp metadirective' was left in either case.
Solution: Check that no 'omp metadirective' is left; additionally, expect
either OMP_NEXT_VARIANT (when offload_nvptx is available) or no 'teams'
directive at all (if not).
gcc/testsuite/ChangeLog:
PR middle-end/118694
* c-c++-common/gomp/attrs-metadirective-3.c: Change to never
expect 'omp metadirective' in the dump. If !offload_nvptx, check
that no 'teams' shows up in the dump; for offload_nvptx, expect
OMP_NEXT_VARIANT and an error about directive between 'target'
and 'teams'.
* c-c++-common/gomp/metadirective-3.c: Likewise.
Tobias Burnus [Fri, 23 May 2025 09:30:48 +0000 (11:30 +0200)]
libgomp.c-c++-common/metadirective-1.c: Expect 'error:' for nvptx compile [PR118694]
OpenMP's 'target teams' is strictly coupled with 'teams'; if the latter
exists, the kernel is launched in directly with multiple teams. Thus,
the host has to know whether the teams construct exists or not. For
#pragma omp target
#pragma omp metadirective when (device={arch("nvptx")}: teams loop)
it is simple when 'nvptx' offloading is not supported, otherwise it depends
on the default device at runtime as the user code asks for a single team for
host fallback and gcn offload and multiple for nvptx offload.
In any case, this commit ensures that no FAIL is printed, whatever a
future solution might look like. Instead of a dg-bogus combined with an
'xfail offload_target_nvptx', one an also argue that a dg-error for
'target offload_target_nvptx' would be more appropriate.
libgomp/ChangeLog:
PR middle-end/118694
* testsuite/libgomp.c-c++-common/metadirective-1.c: xfail when
compiling (also) for nvptx offloading as an error is then expected.
Nathaniel Shead [Thu, 22 May 2025 12:16:22 +0000 (22:16 +1000)]
c++/modules: Fix merge of TLS init functions [PR120363]
The PR notes that we missed setting DECL_CONTEXT on the TLS init
function; we missed this initially because this function is not created
in header units, only named modules.
I also noticed that 'DECL_CONTEXT (fn) = DECL_CONTEXT (var)' was
incorrect: for class members, this ends up having the modules merging
machinery treat the decl as a member function, which breaks when
attempting to dedup against an existing completed class type. Instead
we can just use the global_namespace as the context, because the name of
the function is already mangled appropriately so that we'll match the
correct duplicates.
PR c++/120363
gcc/cp/ChangeLog:
* decl2.cc (get_tls_init_fn): Set context as global_namespace.
(get_tls_wrapper_fn): Likewise.
gcc/testsuite/ChangeLog:
* g++.dg/modules/pr113292_a.H: Move to...
* g++.dg/modules/tls-1_a.H: ...here.
* g++.dg/modules/pr113292_b.C: Move to...
* g++.dg/modules/tls-1_b.C: ...here.
* g++.dg/modules/pr113292_c.C: Move to...
* g++.dg/modules/tls-1_c.C: ...here.
* g++.dg/modules/tls-2_a.C: New test.
* g++.dg/modules/tls-2_b.C: New test.
* g++.dg/modules/tls-2_c.C: New test.
* g++.dg/modules/tls-3.h: New test.
* g++.dg/modules/tls-3_a.H: New test.
* g++.dg/modules/tls-3_b.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com> Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit 66e9a4f3083356b064cc64651edad466a56f762b)
Nathaniel Shead [Fri, 23 May 2025 14:51:49 +0000 (00:51 +1000)]
c++/modules: Fix stream-in of member using-decls [PR120414]
When streaming in a reference to a data member, we have an oversight
where we did not consider USING_DECLs, despite otherwise handling them
here the same as fields. This patch corrects that mistake.
PR c++/120414
gcc/cp/ChangeLog:
* module.cc (trees_in::tree_node): Allow reading a USING_DECL
when streaming tt_data_member.
gcc/testsuite/ChangeLog:
* g++.dg/modules/using-31_a.C: New test.
* g++.dg/modules/using-31_b.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com> Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit 43dddeef7a870ce4db7407f73660504b67a0a919)
Harald Anlauf [Thu, 15 May 2025 19:07:07 +0000 (21:07 +0200)]
Fortran: default-initialization and functions returning derived type [PR85750]
Functions with non-pointer, non-allocatable result and of derived type did
not always get initialized although the type had default-initialization,
and a derived type component had the allocatable or pointer attribute.
Rearrange the logic when to apply default-initialization.
PR fortran/85750
gcc/fortran/ChangeLog:
* resolve.cc (resolve_symbol): Reorder conditions when to apply
default-initializers.
Martin Jambor [Wed, 14 May 2025 10:08:24 +0000 (12:08 +0200)]
tree-sra: Do not create stores into const aggregates (PR111873)
This patch fixes (hopefully the) one remaining place where gimple SRA
was still creating a load into const aggregates. It occurs when there
is a replacement for a load but that replacement is not type
compatible - typically because it is a single field structure.
I have used testcases from duplicates because the original test-case
no longer reproduces for me.
gcc/ChangeLog:
2025-05-13 Martin Jambor <mjambor@suse.cz>
PR tree-optimization/111873
* tree-sra.cc (sra_modify_expr): When processing a load which has
a type-incompatible replacement, do not store the contents of the
replacement into the original aggregate when that aggregate is
const.
gcc/testsuite/ChangeLog:
2025-05-13 Martin Jambor <mjambor@suse.cz>
* gcc.dg/ipa/pr120044-1.c: New test.
* gcc.dg/ipa/pr120044-2.c: Likewise.
* gcc.dg/tree-ssa/pr114864.c: Likewise.
Martin Jambor [Fri, 16 May 2025 15:13:51 +0000 (17:13 +0200)]
ipa: Dump cgraph_node UID instead of order into ipa-clones dump file
Since starting from GCC 15 the order is not unique for any
symtab_nodes but m_uid is, I believe we ought to dump the latter in
the ipa-clones dump, if only so that people can reliably match entries
about new clones to those about removed nodes (if any).
This patch also contains a fixes to a few other places where we have
so far dumped order to our ordinary dumps and which have been
identified by Michal Jires.
gcc/ChangeLog:
2025-05-16 Martin Jambor <mjambor@suse.cz>
* cgraph.h (symtab_node): Make member function get_uid const.
* cgraphclones.cc (dump_callgraph_transformation): Dump m_uid of the
call graph nodes instead of order.
* cgraph.cc (cgraph_node::remove): Likewise.
* ipa-cp.cc (ipcp_lattice<valtype>::print): Likewise.
* ipa-sra.cc (ipa_sra_summarize_function): Likewise.
* symtab.cc (symtab_node::dump_base): Likewise.
Jonathan Wakely [Tue, 20 May 2025 09:53:41 +0000 (10:53 +0100)]
libstdc++: Fix incorrect links to archived SGI STL docs
In r8-7777-g25949ee33201f2 I updated some URLs to point to copies of the
SGI STL docs in the Wayback Machine, because the original pags were no
longer hosted on sgi.com. However, I incorrectly assumed that if one
archived page was at https://web.archive.org/web/20171225062613/... then
all the other pages would be too. Apparently that's not how the Wayback
Machine works, and each page is archived on a different date. That meant
that some of our links were redirecting to archived copies of the
announcement that the SGI STL docs have gone away.
This fixes each URL to refer to a correctly archived copy of the
original docs.
Nathaniel Shead [Sat, 17 May 2025 13:51:07 +0000 (23:51 +1000)]
c++/modules: Fix ICE on merge of instantiation with partial spec [PR120013]
When we import a pending instantiation that matches an existing partial
specialisation, we don't find the slot in the entity map because for
partial specialisations we register the TEMPLATE_DECL but for normal
implicit instantiations we instead register the inner TYPE_DECL.
Because the DECL_MODULE_ENTITY_P flag is set we correctly realise that
it is in the entity map, but ICE when attempting to use that slot in
partition handling.
This patch fixes the issue by detecting this case and instead looking
for the slot for the TEMPLATE_DECL. It doesn't matter that we never add
a slot for the inner decl because we're about to discard it anyway.
PR c++/120013
gcc/cp/ChangeLog:
* module.cc (trees_in::install_entity): Handle re-registering
the inner TYPE_DECL of a partial specialisation.
gcc/testsuite/ChangeLog:
* g++.dg/modules/partial-8.h: New test.
* g++.dg/modules/partial-8_a.C: New test.
* g++.dg/modules/partial-8_b.C: New test.
* g++.dg/modules/partial-8_c.C: New test.
* g++.dg/modules/partial-8_d.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com> Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit b0de7297f2b5670386472229ab795a577c288ecf)
Nathaniel Shead [Mon, 19 May 2025 13:17:16 +0000 (23:17 +1000)]
c++/modules: Always mark tinfo vars as TREE_ADDRESSABLE [PR120350]
We need to mark type info decls as addressable if we take them by
reference; this is done by walking the declaration during parsing and
marking the decl as needed.
However, with modules we don't stream tinfo decls directly; rather we
stream just their name and type and reconstruct them in the importer
directly. This means that any addressable flags are not propagated, and
we error because TREE_ADDRESSABLE is not set despite taking its address.
But tinfo decls should always have TREE_ADDRESSABLE set, as any attempt
to use the tinfo decl will go through build_address anyway. So this
patch fixes the issue by eagerly marking the constructed decl as
TREE_ADDRESSABLE so that modules gets this flag correctly set as well.
PR c++/120350
gcc/cp/ChangeLog:
* rtti.cc (get_tinfo_decl_direct): Mark TREE_ADDRESSABLE.
gcc/testsuite/ChangeLog:
* g++.dg/modules/tinfo-3_a.H: New test.
* g++.dg/modules/tinfo-3_b.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com> Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit 9a6e5a437f0416627ee516f6ef5929cb30c5e498)
Jonathan Wakely [Fri, 16 May 2025 10:54:46 +0000 (11:54 +0100)]
libstdc++: Fix some Clang -Wsystem-headers warnings in <ranges>
libstdc++-v3/ChangeLog:
* include/std/ranges (_ZipTransform::operator()): Remove name of
unused parameter.
(chunk_view::_Iterator, stride_view::_Iterator): Likewise.
(join_with_view): Declare _Iterator and _Sentinel as class
instead of struct.
(repeat_view): Declare _Iterator as class instead of struct.
Jonathan Wakely [Thu, 15 May 2025 18:32:01 +0000 (19:32 +0100)]
libstdc++: Fix std::format of chrono::local_days with {} [PR120293]
Formatting of chrono::local_days with an empty chrono-specs should be
equivalent to inserting it into an ostream, which should use the
overload for inserting chrono::sys_days into an ostream. The
implementation of empty chrono-specs in _M_format_to_ostream takes some
short cuts, and that wasn't being done correctly for chrono::local_days.
libstdc++-v3/ChangeLog:
PR libstdc++/120293
* include/bits/chrono_io.h (_M_format_to_ostream): Add special
case for local_time convertible to local_days.
* testsuite/std/time/clock/local/io.cc: Check formatting of
chrono::local_days.
Jonathan Wakely [Wed, 30 Apr 2025 16:31:01 +0000 (17:31 +0100)]
libstdc++: Fix dangling pointer in fs::path::operator+=(*this) [PR120029]
When concatenating a path we reallocate the left operand's storage to
make room for the new components being added. When the two operands are
the same object, or the right operand is one of the components of the
left operand, the reallocation invalidates the pointers that refer
into the right operand's storage.
The solution in this commit is to detect these aliasing cases and just
do the concatenation in terms of the contained string, as that code
already handles the case where the string aliases the path. The standard
specifies the concatenation in terms of the native() string, so all this
change does is disable the optimized implementation of concatenation for
path objects which attempts to avoid re-parsing the path from the
concatenated string.
The potential loss of performance for this case isn't likely to be an
issue, because concatenating a path with itself (or one of its existing
components) probably isn't a common use case.
The Filesystem TS implementation doesn't have the optimized form of
concatenation and always does it in terms of the native string and
reparsing the whole thing, so doesn't have this bug. A test is added to
confirm that anyway (that test has some slightly different results due
to different behaviour for trailing slashes and implicit "." filenames
in the TS spec).
libstdc++-v3/ChangeLog:
PR libstdc++/120029
* src/c++17/fs_path.cc (path::operator+=(const path&)): Handle
parameters that alias the path or one of its components.
* testsuite/27_io/filesystem/path/concat/120029.cc: New test.
* testsuite/experimental/filesystem/path/concat/120029.cc: New
test.
Jonathan Wakely [Thu, 15 May 2025 10:01:05 +0000 (11:01 +0100)]
libstdc++: Fix std::format_kind primary template for Clang [PR120190]
Although Clang trunk has been adjusted to handle our std::format_kind
definition (because they need to be able to compile the GCC 15.1.0
release), it's probably better to not rely on something that they might
start diagnosing again in future.
Define the primary template in terms of an immediately invoked function
expression, so that we can put a static_assert(false) in the body.
libstdc++-v3/ChangeLog:
PR libstdc++/120190
* include/std/format (format_kind): Adjust primary template to
not depend on itself.
* testsuite/std/format/ranges/format_kind_neg.cc: Adjust
expected errors. Check more invalid specializations.
Reviewed-by: Tomasz Kamiński <tkaminsk@redhat.com> Reviewed-by: Daniel Krügler <daniel.kruegler@gmail.com>
(cherry picked from commit c65725eccbabf3b9b5965f27fff2d3b9f6c75930)
Tobias Burnus [Thu, 15 May 2025 07:15:21 +0000 (09:15 +0200)]
OpenMP/Fortran: Fix allocatable-component mapping of derived-type array comps
The check whether the location expression in map clause has allocatable
components was failing for some derived-type array expressions such as
map(var%tiles(1))
as the compiler produced
_4 = var.tiles;
MEMREF(_4, _5);
This commit now also handles this case.
gcc/fortran/ChangeLog:
* trans-openmp.cc (gfc_omp_deep_mapping_do): Handle SSA_NAME if
a def_stmt is available.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/alloc-comp-4.f90: New test.
Tobias Burnus [Wed, 14 May 2025 18:06:49 +0000 (20:06 +0200)]
OpenMP: Fix mapping of zero-sized arrays with non-literal size: map(var[:n]), n = 0
For map(ptr[:0]), the used map kind is GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
and it is permitted that 'ptr' does not exist. 'ptr' is set to the device
pointee if it exists or to the host value otherwise.
For map(ptr[:3]), the variable is first mapped and then ptr is updated to point
to the just-mapped device data; the attachment uses GOMP_MAP_ATTACH.
For map(ptr[:n]), generates always a GOMP_MAP_ATTACH, but when n == 0, it
was failing with:
"pointer target not mapped for attach"
The solution is not to fail but first to check whether it was mapped before.
It turned out that for the mapping part, GCC adds a run-time check whether
n == 0 - and uses GOMP_MAP_ZERO_LEN_ARRAY_SECTION for the mapping.
Thus, we just have to check whether there such a mapping for the address
for which the GOMP_MAP_ATTACH. was requested. And, if there was, the
error diagnostic can be skipped.
Unsurprisingly, this issue occurs in real-world code; it was detected in
a code that distributes work via MPI and for some processes, some bounds
ended up to be zero.
libgomp/ChangeLog:
* target.c (gomp_attach_pointer): Return bool; accept additional
bool to optionally silence the fatal pointee-not-found error.
(gomp_map_vars_internal): If the pointee could not be found,
check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
* libgomp.h (gomp_attach_pointer): Update prototype.
* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
calls.
* testsuite/libgomp.c/target-map-zero-sized.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.
While the tests checked whether the CUDA/HIP runtime is available
before processing them, the execution was then done unconditionally,
leading to FAIL when the default device was the host (or the wrong
offload device).
Now the test is only executed ('run') when the default device is an
Nvidia or AMD GPU (depending on the test case, cf. the test file name).
Otherwise, only a 'link' test is done. (Except when the effective-target
check cannot find the runtime lib - then the test is skipped [as before].)
Note: The cublas/hipblas tests use variant functions and iterate over
all devices, such that the cublas or hipblas, respectively, is only
called when the active device is an AMD or Nvidia device, respectively,
while for the host and other device types the fallback is called.
libgomp/ChangeLog:
* testsuite/libgomp.c/interop-cuda-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_nvptx".
* testsuite/libgomp.c/interop-cuda-libonly.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: Likewise.
* testsuite/libgomp.c/interop-hip-amd-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_gcn".
* testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: Likewise.
When host memory is device accessible - independent whether mapping is done or
not (i.e. self map), the 'vtab' pointer becomes accessible, which stores the
dynamic type's type and size information.
In principle, we want to test: USM available but mapping is still done, but
as there is no simple + reliable not-crashing way to test for this, those
checks are skipped in the (pre)existing test file map-alloc-comp-9.f90.
Or rather: those are only active with self-maps, which is currently only true
for the host.
This commit adds map-alloc-comp-9-usm.f90 which runs the same test with
'omp requires unified_shared_memory'. While OpenMP permits both actual
mapping and self maps with this flag, it in theory covers the missing cases.
However, currently, GCC always uses self maps with USM. Still, having a
device-run self-maps check is better than nothing, even if it misses the
most interesting case.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/map-alloc-comp-9.f90: Process differently
when USE_USM_REQUIREMENT is set.
* testsuite/libgomp.fortran/map-alloc-comp-9-usm.f90: New test.
Thomas Schwinge [Mon, 5 May 2025 08:19:30 +0000 (10:19 +0200)]
'libgomp.c/interop-hsa.c': GCN offloading only
Fix-up for commit 8d84ea28510054fbbb8a2b7441916bd75e29163f
"OpenMP, GCN: Add interop-hsa testcase", which added 'libgomp.c/interop-hsa.c'.
If nvptx offloading compilation is enabled in addition to GCN, the former ICEs:
Tobias Burnus [Thu, 1 May 2025 15:39:42 +0000 (15:39 +0000)]
OpenMP: Restore lost Fortran testcase for 'omp allocate'
This testcase, which is present on the OG13 and OG14 branches, was
overlooked when the Fortran support for 'omp allocate' was added to
mainline (commit d4b6d147920b93297e621124a99ed01e7e310d92 from
December 2023).
libgomp/ChangeLog
* testsuite/libgomp.fortran/allocate-8a.f90: New test.
libgomp/testsuite: Fix hip_header_nvidia check, add workaround to test
This is all about using the AMD's HIP header files with
__HIP_PLATFORM_NVIDIA__ defined, i.e. HIP with Nvidia/CUDA; in that case,
HIP is a thin layer on top of CUDA.
First, the check_effective_target_gomp_hip_header_nvidia check failed;
to fix it, -Wno-deprecated-declarations was added - and likewise to the
two affected testcases that actually used the HIP headers on Nvidia.
Doing so, the HIP tested was successful but the HIP-BLAS one showed two
issues:
* One seems to be related to include search paths as the HIP header uses
#include "library_types.h" to include that CUDA header. Seemingly, it
tried to included (again) the HIP header hip/library_types.h, not the
CUDA one. I guess, some tweaking of -isystem vs. -I could have
prevented this, but the simpler workaround was to just explicitly
include the CUDA one before the HIP header files.
* Once done, everything compiled but linking failed as the association
between three HIP-BLAS functions and their CUDA-BLAS ones did not
work. Solution: Just add three #define for mapping them.
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp
(check_effective_target_gomp_hip_header_nvidia): Compile with
"-Wno-deprecated-declarations".
* testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise.
* testsuite/libgomp.c/interop-hipblas-nvidia-full.c: Likewise.
* testsuite/libgomp.c/interop-hipblas.h: Add workarounds
when using the HIP headers with __HIP_PLATFORM_NVIDIA__.
Add checks for nowait/depend and for checks that the returned
CUDA, CUDA_DRIVER and HIP interop objects actually work.
While the CUDA/CUDA_DRIVER ones are only for Nvidia GPUs, HIP
works on both AMD and Nvidia GPUs; on Nvidia GPUs, it is a
very thin wrapper around CUDA.
For Fortran, only a HIP test has been added - using hipfort.
While libgomp.c-c++-common/interop-2.c always works - even without
GPU - and checks for depend / nowait, all others require that
runtime libraries are found at link (and execution) time:
For Nvidia GPUs, libcuda + libcudart or libcublas,
For AMD GPUs, libamdhip64 or libhipblas.
The header files and hipfort modules do not need to be present as a
fallback has been implemented, but if they are, they get used.
Due to the combinations, the basic 1x C/C++, 4x C and 1x Fortran tests
yield 1x C/C++, 14x C and 4 Fortran run-test files.
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp (check_effective_target_openacc_cublas,
check_effective_target_openacc_cudart): Update description as
the check requires more.
(check_effective_target_openacc_libcuda,
check_effective_target_openacc_libcublas,
check_effective_target_openacc_libcudart,
check_effective_target_gomp_hip_header_amd,
check_effective_target_gomp_hip_header_nvidia,
check_effective_target_gomp_hipfort_module,
check_effective_target_gomp_libamdhip64,
check_effective_target_gomp_libhipblas): New.
* testsuite/libgomp.c-c++-common/interop-2.c: New test.
* testsuite/libgomp.c/interop-cublas-full.c: New test.
* testsuite/libgomp.c/interop-cublas-libonly.c: New test.
* testsuite/libgomp.c/interop-cuda-full.c: New test.
* testsuite/libgomp.c/interop-cuda-libonly.c: New test.
* testsuite/libgomp.c/interop-hip-amd-full.c: New test.
* testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: New test.
* testsuite/libgomp.c/interop-hip-nvidia-full.c: New test.
* testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: New test.
* testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: New test.
* testsuite/libgomp.c/interop-hip.h: New test.
* testsuite/libgomp.c/interop-hipblas-amd-full.c: New test.
* testsuite/libgomp.c/interop-hipblas-amd-no-hip-header.c: New test.
* testsuite/libgomp.c/interop-hipblas-nvidia-full.c: New test.
* testsuite/libgomp.c/interop-hipblas-nvidia-no-headers.c: New test.
* testsuite/libgomp.c/interop-hipblas-nvidia-no-hip-header.c: New test.
* testsuite/libgomp.c/interop-hipblas.h: New test.
* testsuite/libgomp.fortran/interop-hip-amd-full.F90: New test.
* testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: New test.
* testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: New test.
* testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: New test.
* testsuite/libgomp.fortran/interop-hip.h: New test.
Alpha: Fix base block alignment calculation regression
In determination of base block alignment we only examine a COMPONENT_REF
tree node at hand without ever checking if its ultimate alignment has
been reduced by the combined offset going back to the outermost object.
Consequently cases have been observed where quadword accesses have been
produced for a memory location referring a nested struct member only
aligned to the longword boundary, causing emulation to trigger.
Address this issue by recursing into COMPONENT_REF tree nodes until the
outermost one has been reached, which is supposed to be a MEM_REF one,
accumulating the offset as we go, fixing a commit e0dae4da4c45 ("Alpha:
Also use tree information to get base block alignment") regression.
Bail out and refrain from using tree information for alignment if we end
up at something different or we are unable to calculate the offset at
any point.
gcc/
* config/alpha/alpha.cc
(alpha_get_mem_rtx_alignment_and_offset): Recurse into
COMPONENT_REF nodes.
gcc/testsuite/
* gcc.target/alpha/memcpy-nested-offset-long.c: New file.
* gcc.target/alpha/memcpy-nested-offset-quad.c: New file.
projects / thirdparty / gcc.git / log
