aboutsummaryrefslogtreecommitdiff
path: root/gcc
AgeCommit message (Collapse)AuthorFilesLines
2022-07-05OpenMP: Handle tofrom with target enter/exit dataTobias Burnus12-10/+166
In 5.2, a map clause can be map-entering or map-exiting, either containing 'tofrom'. The main reason for this is permit 'map(x)' with 'omp target enter/exit data', avoiding to specify 'to:/from:' explicitly. (OpenMP defaults to 'tofrom'.) gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Accept tofrom map-type modifier but use 'to' / 'from' internally. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data): Accept tofrom map-type modifier but use 'to' / 'from' internally. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_namelist): For the map-type, also handle the always modifer and release/delete. * openmp.cc (resolve_omp_clauses): Accept tofrom map-type modifier for target enter/exit data, but use 'to' / 'from' internally. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.2): Mark target enter/exit data with fromto as implemented. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-data-2.c: New test. * c-c++-common/gomp/target-data-3.c: New test. * gfortran.dg/gomp/target-data-1.f90: New test. * gfortran.dg/gomp/target-data-2.f90: New test. (cherry picked from commit 9a668532fb19e7c57aa595a26ce3f0d95f9cbb1b)
2022-07-05gcn: Remove useless register keywordTobias Burnus2-1/+9
gcc/ChangeLog: * config/gcn/gcn-protos.h (print_operand_address): Remove register keyword on 'rtx addr' argument. (cherry picked from commit 7780dc5b2d02785186583fc8eced3c9e3aec4552)
2022-07-05amdgcn: test global constructorsAndrew Stubbs2-1/+8
The tests are disabled for historical reasons only. gcc/testsuite/ChangeLog: * lib/target-supports.exp (check_effective_target_global_constructor): Remove amdgcn. (cherry picked from commit a1f8a3860fe5c83a023688c29636b5abe03db949)
2022-07-05amdgcn: remove obsolete assembler workaroundsAndrew Stubbs2-29/+16
This nonsense is no longer required, now that the minimum supported assembler version is LLVM 13.0.1. gcc/ChangeLog: * config/gcn/gcn.md (*movbi): Remove assembler bug workarounds. (jump): Likewise. (movdi_symbol_save_scc): Likewise. (cherry picked from commit b06a282921c71bbc5cab69bc515804bd80f55e92)
2022-07-05openmp: Conforming device numbers and omp_{initial,invalid}_deviceJakub Jelinek2-10/+73
OpenMP 5.2 changed once more what device numbers are allowed. In 5.1, valid device numbers were [0, omp_get_num_devices()]. 5.2 makes also -1 valid (calls it omp_initial_device), which is equivalent in behavior to omp_get_num_devices() number but has the advantage that it is a constant. And it also introduces omp_invalid_device which is also a constant with implementation defined value < -1. That value should act like sNaN, any time any device construct (GOMP_target*) or OpenMP runtime API routine is asked for such a device, the program is terminated. And if OMP_TARGET_OFFLOAD=mandatory, all non-conforming device numbers (which is all but [-1, omp_get_num_devices()] other than omp_invalid_device) must be treated like omp_invalid_device. For device constructs, we have a compatibility problem, we've historically used 2 magic negative values to mean something special. GOMP_DEVICE_ICV (-1) means device clause wasn't present, pick the omp_get_default_device () number GOMP_DEVICE_FALLBACK (-2) means the host device (this is used e.g. for #pragma omp target if (cond) where if cond is false, we pass -2 But 5.2 requires that omp_initial_device is -1 (there were discussions about it, advantage of -1 is that one can say iterate over the [-1, omp_get_num_devices()-1] range to get all devices starting with the host/initial one. And also, if user passes -2, unless it is omp_invalid_device, we need to treat it like non-conforming with OMP_TARGET_OFFLOAD=mandatory. So, the patch does on the compiler side some number remapping, user_device_num >= -2U ? user_device_num - 1 : user_device_num. This remapping is done at compile time if device clause has constant argument, otherwise at runtime, and means that for user -1 (omp_initial_device) we pass -2 to GOMP_* in the runtime library where it treats it like host fallback, while -2 is remapped to -3 (one of the non-conforming device numbers, for those it doesn't matter which one is which). omp_invalid_device is then -4. For the OpenMP device runtime APIs, no remapping is done. This patch doesn't deal with the initial default-device-var for OMP_TARGET_OFFLOAD=mandatory , the spec says that the inital ICV value for that should in that case depend on whether there are any offloading devices or not (if not, should be omp_invalid_device), but that means we can't determine the number of devices lazily (and let libraries have the possibility to register their offloading data etc.). 2022-06-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-expand.cc (expand_omp_target): Remap user provided device clause arguments, -1 to -2 and -2 to -3, either at compile time if constant, or at runtime. include/ * gomp-constants.h (GOMP_DEVICE_INVALID): Define. libgomp/ * omp.h.in (omp_initial_device, omp_invalid_device): New enumerators. * omp_lib.f90.in (omp_initial_device, omp_invalid_device): New parameters. * omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise. * target.c (resolve_device): Add remapped argument, handle GOMP_DEVICE_ICV only if remapped is true (and clear remapped), for negative values, treat GOMP_DEVICE_FALLBACK as fallback only if remapped, otherwise treat omp_initial_device that way. For omp_invalid_device, always emit gomp_fatal, even when OMP_TARGET_OFFLOAD isn't mandatory. (GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext, GOMP_target_update, GOMP_target_update_ext, GOMP_target_enter_exit_data): Pass true as remapped argument to resolve_device. (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy_check, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_get_mapped_ptr, omp_target_is_accessible): Pass false as remapped argument to resolve_device. Treat omp_initial_device the same as gomp_get_num_devices (). Don't bypass resolve_device calls if device_num is negative. (omp_pause_resource): Treat omp_initial_device the same as gomp_get_num_devices (). Call resolve_device. * icv-device.c (omp_set_default_device): Always set to device_num even when it is negative. * libgomp.texi: Document that Conforming device numbers, omp_initial_device and omp_invalid_device is implemented. * testsuite/libgomp.c/target-41.c (main): Add test with omp_initial_device. * testsuite/libgomp.c/target-45.c: New test. * testsuite/libgomp.c/target-46.c: New test. * testsuite/libgomp.c/target-47.c: New test. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add test with omp_initial_device. Use -5 instead of -1 for negative value test. * testsuite/libgomp.fortran/target-is-accessible-1.f90 (main): Likewise. Reorder stop numbers. (cherry picked from commit 1158fe43407568f20415b16575ddbfff216bf8b6)
2022-07-05OpenMP: Handle ancestor:1 with discover_declare_targetTobias Burnus5-10/+135
gcc/ * omp-offload.cc (omp_discover_declare_target_tgt_fn_r, omp_discover_declare_target_fn_r): Don't walk reverse-offload target regions. gcc/testsuite/ * c-c++-common/gomp/reverse-offload-1.c: New. (cherry picked from commit 209de00fdb9da90a727337d6e752ea583a11d982)
2022-07-05OpenMP: Fortran - fix ancestor's requires reverse_offload checkTobias Burnus2-1/+77
gcc/fortran/ * openmp.cc (gfc_match_omp_clauses): Check also parent namespace for 'requires reverse_offload'. gcc/testsuite/ * gfortran.dg/gomp/target-device-ancestor-5.f90: New test. (cherry picked from commit 5e5deac508e3025e2d2c36212aa52d52001b893d)
2022-07-05openmp: Add support for OpenMP 5.2 linear clause syntax for C/C++Jakub Jelinek17-37/+508
The syntax for linear clause changed in 5.2, the original syntax which is still valid is: linear (var1, var2) linear (var3, var4 : step1) The 4.5 syntax with modifiers like: linear (val (var5, var6)) linear (val (var7, var8) : step2) is still supported in 5.2, but is deprecated there. Instead, one can use a new syntax: linear (var9, var10 : val) linear (var11, var12 : step (step3), val) As val, ref, uval or step (someexpr) can be valid expressions (and especially in C++ can be const / constexpr / consteval), the spec says that when the whole step expression is val (or ref or uval) or step ( ... ) then it is the new modifier syntax, one can use + 0 or 0 + or 1 * or * 1 or ()s to say it is the old step expression. Also, 5.2 now allows val modifier to be specified even outside of declare simd (but not the other modifiers). I've implemented this for the new modifier syntax only, the old one keeps the old restriction (which is why OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER flag has been introduced). 2022-06-07 Jakub Jelinek <jakub@redhat.com> gcc/ * tree.h (OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER): Define. * tree-pretty-print.cc (dump_omp_clause) <case OMP_CLAUSE_LINEAR>: Adjust clause printing style depending on OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER. gcc/c/ * c-parser.cc (c_parser_omp_clause_linear): Parse OpenMP 5.2 style linear clause modifiers. Set OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER flag on the clauses when old style modifiers are used. * c-typeck.cc (c_finish_omp_clauses): Only reject linear clause with val modifier on simd or for if the old style modifiers are used. gcc/cp/ * parser.cc (cp_parser_omp_clause_linear): Parse OpenMP 5.2 style linear clause modifiers. Set OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER flag on the clauses when old style modifiers are used. * semantics.cc (finish_omp_clauses): Only reject linear clause with val modifier on simd or for if the old style modifiers are used. gcc/fortran/ * trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_LINEAR_OLD_LINEAR_MODIFIER on OMP_CLAUSE_LINEAR clauses unconditionally for now. gcc/testsuite/ * c-c++-common/gomp/linear-2.c: New test. * c-c++-common/gomp/linear-3.c: New test. * g++.dg/gomp/linear-3.C: New test. * g++.dg/gomp/linear-4.C: New test. * g++.dg/gomp/linear-5.C: New test. (cherry picked from commit 03b71406323ddc065b1d7837d8b43b17e4b048b5)
2022-07-05arm: reinstate HAVE_GAS_ARM_EXTENDED_ARCHAndrew Stubbs4-0/+67
The check was removed by accident. gcc/ChangeLog: * config.in: Regenerate. * configure: Regenerate. * configure.ac: Reinstate HAVE_GAS_ARM_EXTENDED_ARCH test. (cherry picked from commit 36bd6eafb6062f1fb92a994538c6ed017ced670b)
2022-07-05gfortran.dg/gomp/scope-6.f90: Add \\ to scan-tree-dumpTobias Burnus2-1/+8
Missed git add for the hot fix before committing r13-982-gff35a75473d28205e52ecbcf9e6b5107b8b5ab90 gcc/testsuite/ * gfortran.dg/gomp/scope-6.f90: Fix dg-final scan-tree-dump. (cherry picked from commit 6a098f4e16351617ebbb9ad66fb802e8c4a57c77)
2022-07-05OpenMP/Fortran: Add support for firstprivate and allocate clauses on scope ↵Tobias Burnus5-1/+49
construct Fortran commit to C/C++/backend commit r13-862-gf38b20d68fade5a922b9f68c4c3841e653d1b83c gcc/fortran/ChangeLog: * openmp.cc (OMP_SCOPE_CLAUSES): Add firstprivate and allocate. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.2): Mark scope w/ firstprivate/allocate as Y. * testsuite/libgomp.fortran/scope-2.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/scope-5.f90: New test. * gfortran.dg/gomp/scope-6.f90: New test. (cherry picked from commit ff35a75473d28205e52ecbcf9e6b5107b8b5ab90)
2022-07-05openmp: Add support for firstprivate and allocate clauses on scope constructJakub Jelinek11-5/+91
OpenMP 5.2 adds support for firstprivate and allocate clauses on the scope construct and this patch adds that support to GCC. 5.2 unfortunately (IMNSHO mistakenly) marked scope construct as worksharing, which implies that it isn't possible to nest inside of it other scope, worksharing loop, sections, explicit barriers, single etc. which would make scope far less useful. I'm not implementing that part, keeping the 5.1 behavior here, and will file an issue to revert that for OpenMP 6.0. But, for firstprivate it keeps the restriction that is now implied from worksharing construct that listed var can't be private in outer context, where for reduction 5.1 had similar restriction explicit even for scope and 5.2 has it implicitly through worksharing construct. 2022-05-31 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.cc (build_outer_var_ref): For code == OMP_CLAUSE_ALLOCATE allow var to be private in the outer context. (lower_private_allocate): Pass OMP_CLAUSE_ALLOCATE as last argument to build_outer_var_ref. gcc/c/ * c-parser.cc (OMP_SCOPE_CLAUSE_MASK): Add firstprivate and allocate clauses. gcc/cp/ * parser.cc (OMP_SCOPE_CLAUSE_MASK): Add firstprivate and allocate clauses. gcc/testsuite/ * c-c++-common/gomp/scope-5.c: New test. * c-c++-common/gomp/scope-6.c: New test. * g++.dg/gomp/attrs-1.C (bar): Add firstprivate and allocate clauses to scope construct. * g++.dg/gomp/attrs-2.C (bar): Likewise. libgomp/ * testsuite/libgomp.c-c++-common/allocate-1.c (foo): Add testcase for scope construct with allocate clause. * testsuite/libgomp.c-c++-common/allocate-3.c (foo): Likewise. * testsuite/libgomp.c-c++-common/scope-2.c: New test. (cherry picked from commit f38b20d68fade5a922b9f68c4c3841e653d1b83c)
2022-07-05OpenMP/Fortran: Add support for enter clause on declare targetTobias Burnus7-28/+86
Fortran version to C/C++ commit r13-797-g0ccba4ed8571c18c7015413441e971 gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_clauses): Handle OMP_LIST_ENTER. * gfortran.h: Add OMP_LIST_ENTER. * openmp.cc (enum omp_mask2, OMP_DECLARE_TARGET_CLAUSES): Add OMP_CLAUSE_ENTER. (gfc_match_omp_clauses, gfc_match_omp_declare_target, resolve_omp_clauses): Handle 'enter' clause. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.2): Mark 'enter' clause as supported. * testsuite/libgomp.fortran/declare-target-1.f90: Extend to test explicit 'to' and 'enter' clause. * testsuite/libgomp.fortran/declare-target-2.f90: Update accordingly. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-target-2.f90: Add 'enter' clause test. * gfortran.dg/gomp/declare-target-4.f90: Likewise. (cherry picked from commit e3803f9cbb67efef7462589fe50d00b404356bc9)
2022-07-05openmp: Adjust diagnostics about same variable in link and to/enter clausesJakub Jelinek6-7/+60
On Fri, May 27, 2022 at 04:52:17PM +0200, Tobias Burnus wrote: > The 'TO'/'ENTER' usage is first stored in a linked list – and > then as attribute to the symbol. I am not sure how to handle it best. This reminds me I've left the C/C++ FE diagnostics about mixing link and to/enter on the same variable in separate directives as is, so it always talked about mixing link and to clauses. This patch adjusts it, so that if link is first, it talks about the clause actually used and if link is later, uses to or enter together in the wording. 2022-05-28 Jakub Jelinek <jakub@redhat.com> gcc/c/ * c-parser.cc (c_parser_omp_declare_target): If OMP_CLAUSE_LINK was seen first, use "%<to%>" or "%<enter%>" depending on OMP_CLAUSE_ENTER_TO of the current clause, otherwise use "%<to%> or %<enter%>" wording. gcc/cp/ * parser.cc (handle_omp_declare_target_clause): If OMP_CLAUSE_LINK was seen first, use "%<to%>" or "%<enter%>" depending on OMP_CLAUSE_ENTER_TO of the current clause, otherwise use "%<to%> or %<enter%>" wording. gcc/testsuite/ * c-c++-common/gomp/declare-target-2.c: Add further tests for mixing of link and to/enter clauses on separate directives. (cherry picked from commit 931249276acb4bdab778fe2bf4e81a1a31a78e6c)
2022-07-05openmp: Add support for enter clause on declare targetJakub Jelinek21-42/+171
OpenMP 5.1 and earlier had 2 different uses of to clause, one for target update construct with one semantics, and one for declare target directive with a different semantics. Under the hood we were using OMP_CLAUSE_TO_DECLARE to represent the latter. OpenMP 5.2 renamed the declare target clause to to enter, the old one is kept as a deprecated alias. As we are far from having full OpenMP 5.2 support, this patch adds support for the enter clause (and renames OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER with a flag to tell the spelling of the clause for better diagnostics), but doesn't deprecate the to clause on declare target just yet (that should be done as one of the last steps in 5.2 support). 2022-05-27 Jakub Jelinek <jakub@redhat.com> gcc/ * tree-core.h (enum omp_clause_code): Rename OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER. * tree.h (OMP_CLAUSE_ENTER_TO): Define. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Rename OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, if OMP_CLAUSE_ENTER_TO, print "to" instead of "enter". * tree-nested.cc (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_ENTER. gcc/c/ * c-parser.cc (c_parser_omp_clause_name): Parse enter clause. (c_parser_omp_all_clauses): For to clause on declare target, use OMP_CLAUSE_ENTER clause with OMP_CLAUSE_ENTER_TO instead of OMP_CLAUSE_TO_DECLARE clause. Handle PRAGMA_OMP_CLAUSE_ENTER. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add enter clause. (c_parser_omp_declare_target): Use OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. * c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, to OMP_CLAUSE_ENTER_TO use "to" as clause name in diagnostics instead of omp_clause_code_name[OMP_CLAUSE_CODE (c)]. gcc/cp/ * parser.cc (cp_parser_omp_clause_name): Parse enter clause. (cp_parser_omp_all_clauses): For to clause on declare target, use OMP_CLAUSE_ENTER clause with OMP_CLAUSE_ENTER_TO instead of OMP_CLAUSE_TO_DECLARE clause. Handle PRAGMA_OMP_CLAUSE_ENTER. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add enter clause. (cp_parser_omp_declare_target): Use OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. * semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, to OMP_CLAUSE_ENTER_TO use "to" as clause name in diagnostics instead of omp_clause_code_name[OMP_CLAUSE_CODE (c)]. gcc/testsuite/ * c-c++-common/gomp/clauses-3.c: Add tests with enter clause instead of to or modify some existing to clauses to enter. * c-c++-common/gomp/declare-target-1.c: Likewise. * c-c++-common/gomp/declare-target-2.c: Likewise. * c-c++-common/gomp/declare-target-3.c: Likewise. * g++.dg/gomp/attrs-9.C: Likewise. * g++.dg/gomp/declare-target-1.C: Likewise. libgomp/ * testsuite/libgomp.c-c++-common/target-40.c: Modify some existing to clauses to enter. * testsuite/libgomp.c/target-41.c: Likewise. (cherry picked from commit 0ccba4ed8571c18c7015413441e971d4863c3644)
2022-07-05GCN: Add gfx908/gfx90a to -march/-mtune in invoke.texiTobias Burnus2-1/+13
gcc/ * doc/invoke.texi (AMD GCN Options): Add gfx908/gfx90a. (cherry picked from commit 2a790686fd11bc90032ff67c996c5a2752305625)
2022-07-05OpenMP: Support nowait with Fortran [PR105378]Tobias Burnus4-1/+32
Fortran part to C/C++/libgomp commit r13-724-gb43836914bdc2a37563cf31359b2c4803bfe4374 gcc/fortran/ PR c/105378 * openmp.cc (gfc_match_omp_taskwait): Accept nowait. gcc/testsuite/ PR c/105378 * gfortran.dg/gomp/taskwait-depend-nowait-1.f90: New. libgomp/ PR c/105378 * libgomp.texi (OpenMP 5.1): Set 'taskwait nowait' to 'Y'. * testsuite/libgomp.fortran/taskwait-depend-nowait-1.f90: New. (cherry picked from commit 4fb2b4f7ea6b80ae75d3efb6f86e7c6179080535)
2022-07-05openmp: Add taskwait nowait depend support [PR105378]Jakub Jelinek10-12/+91
This patch adds support for (so far C/C++) #pragma omp taskwait nowait depend(...) directive, which is like #pragma omp task depend(...) ; but slightly optimized on the library side, so that it creates the task only for the purpose of dependency tracking and doesn't actually schedule it and wait for it when the dependencies are satisfied, instead makes its dependencies satisfied right away. 2022-05-24 Jakub Jelinek <jakub@redhat.com> PR c/105378 gcc/ * omp-builtins.def (BUILT_IN_GOMP_TASKWAIT_DEPEND_NOWAIT): New builtin. * gimplify.cc (gimplify_omp_task): Diagnose taskwait with nowait clause but no depend clauses. * omp-expand.cc (expand_taskwait_call): Use BUILT_IN_GOMP_TASKWAIT_DEPEND_NOWAIT rather than BUILT_IN_GOMP_TASKWAIT_DEPEND if nowait clause is present. gcc/c/ * c-parser.cc (OMP_TASKWAIT_CLAUSE_MASK): Add nowait clause. gcc/cp/ * parser.cc (OMP_TASKWAIT_CLAUSE_MASK): Add nowait clause. gcc/testsuite/ * c-c++-common/gomp/taskwait-depend-nowait-1.c: New test. libgomp/ * libgomp_g.h (GOMP_taskwait_depend_nowait): Declare. * libgomp.map (GOMP_taskwait_depend_nowait): Export at GOMP_5.1.1. * task.c (empty_task): New function. (gomp_task_run_post_handle_depend_hash): Declare earlier. (gomp_task_run_post_handle_depend): Declare. (GOMP_task): Optimize fn == empty_task if there is nothing to wait for. (gomp_task_run_post_handle_dependers): Optimize task->fn == empty_task. (GOMP_taskwait_depend_nowait): New function. * testsuite/libgomp.c-c++-common/taskwait-depend-nowait-1.c: New test. (cherry picked from commit b43836914bdc2a37563cf31359b2c4803bfe4374)
2022-07-05libgomp: Add new runtime routines omp_target_memcpy_async and ↵Marcel Vollweiler2-0/+10
omp_target_memcpy_rect_async This patch adds two new OpenMP runtime routines: omp_target_memcpy_async and omp_target_memcpy_rect_async. Both functions are introduced in OpenMP 5.1 as asynchronous variants of omp_target_memcpy and omp_target_memcpy_rect. In contrast to the synchronous variants, the asynchronous functions have two additional function parameters to allow the specification of task dependences: int depobj_count omp_depend_t *depobj_list integer(c_int), value :: depobj_count integer(omp_depend_kind), optional :: depobj_list(*) The implementation splits the synchronous functions into two parts: (a) check and (b) copy. Then (a) is used in the asynchronous functions for the sequential part, and the actual copy process (b) is executed in a new created task. The sequential part (a) takes into account the requirements for the return values: "The routine returns zero if successful. Otherwise, it returns a non-zero value." (omp_target_memcpy_async, OpenMP 5.1 spec, section 3.8.7) "An application can determine the number of inclusive dimensions supported by an implementation by passing NULL pointers (or C_NULL_PTR, for Fortran) for both dst and src. The routine returns the number of dimensions supported by the implementation for the specified device numbers. No copy operation is performed." (omp_target_memcpy_rect_async, OpenMP 5.1 spec, section 3.8.8) Due to asynchronicity an error is thrown if the asynchronous memcpy is not successful (in contrast to the synchronous functions which use a return value unequal to zero). gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added target_memcpy_async and target_memcpy_rect_async to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_target_memcpy_async and omp_target_memcpy_rect_async. * libgomp.texi: Both functions are now supported. * omp.h.in: Added omp_target_memcpy_async and omp_target_memcpy_rect_async. * omp_lib.f90.in: Added interfaces for both new functions. * omp_lib.h.in: Likewise. * target.c (ialias_redirect): Added for GOMP_task. (omp_target_memcpy): Restructured into check and copy part. (omp_target_memcpy_check): New helper function for omp_target_memcpy and omp_target_memcpy_async that checks requirements. (omp_target_memcpy_copy): New helper function for omp_target_memcpy and omp_target_memcpy_async that performs the memcpy. (omp_target_memcpy_async_helper): New helper function that is used in omp_target_memcpy_async for the asynchronous task. (omp_target_memcpy_async): Added. (omp_target_memcpy_rect): Restructured into check and copy part. (omp_target_memcpy_rect_check): New helper function for omp_target_memcpy_rect and omp_target_memcpy_rect_async that checks requirements. (omp_target_memcpy_rect_copy): New helper function for omp_target_memcpy_rect and omp_target_memcpy_rect_async that performs the memcpy. (omp_target_memcpy_rect_async_helper): New helper function that is used in omp_target_memcpy_rect_async for the asynchronous task. (omp_target_memcpy_rect_async): Added. * task.c (ialias): Added for GOMP_task. * testsuite/libgomp.c-c++-common/target-memcpy-async-1.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-async-2.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-rect-async-1.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-rect-async-2.c: New test. * testsuite/libgomp.fortran/target-memcpy-async-1.f90: New test. * testsuite/libgomp.fortran/target-memcpy-async-2.f90: New test. * testsuite/libgomp.fortran/target-memcpy-rect-async-1.f90: New test. * testsuite/libgomp.fortran/target-memcpy-rect-async-2.f90: New test. (cherry picked from commit 6c420193e86b39a09304b2845335571eefe24d5d)
2022-07-05OpenMP: Add Fortran support for inoutset depend-kindTobias Burnus10-8/+51
Fortran additions to the C/C++ + ME/libgomp commit r13-556-g2c16eb3157f86ae561468c540caf8eb326106b5f gcc/fortran/ChangeLog: * gfortran.h (enum gfc_omp_depend_op): Add OMP_DEPEND_INOUTSET. (gfc_omp_clauses): Enlarge ENUM_BITFIELD. * dump-parse-tree.cc (show_omp_namelist, show_omp_clauses): Handle 'inoutset' depend modifier. * openmp.cc (gfc_match_omp_clauses, gfc_match_omp_depobj): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses, gfc_trans_omp_depobj): Likewise. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1): Set 'inoutset' to Y. (OpenMP Context Selectors): Add missing comma. * testsuite/libgomp.fortran/depend-5.f90: Add inoutset test. * testsuite/libgomp.fortran/depend-6.f90: Likewise. * testsuite/libgomp.fortran/depend-7.f90: Likewise. * testsuite/libgomp.fortran/depend-inoutset-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/all-memory-1.f90: Add inoutset test. * gfortran.dg/gomp/all-memory-2.f90: Likewise. * gfortran.dg/gomp/depobj-1.f90: Likewise. * gfortran.dg/gomp/depobj-2.f90: Likewise. (cherry picked from commit ba8563693fb4eb1317fcbd54883516c3873782f4)
2022-07-05OpenMP: Skip target-nesting warning for reverse offloadTobias Burnus4-0/+53
gcc/ChangeLog: * omp-low.cc (check_omp_nesting_restrictions): Skip warning for target inside target if inner is reverse offload. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-device-ancestor-5.c: New test. (cherry picked from commit 47554478a13f64bff1ee4b9bb0319ae63d42ca52)
2022-07-05Daily bump.GCC Administrator4-1/+81
2022-07-04gcn/t-omp-device: Add 'amdgcn' as 'arch' [PR105602]Tobias Burnus3-2/+11
Improve cross-compiler handling. gcc/ChangeLog: PR target/105602 * config/gcn/t-omp-device (arch): Add 'amdgcn' besides existing 'gcn'. * config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Likewise. (cherry picked from commit ebe5dace8c318b38f42cfe1d148e90dcdfadb31e)
2022-07-04openmp: Add support for inoutset depend-kindJakub Jelinek17-27/+211
This patch adds support for inoutset depend-kind in depend clauses. It is very similar to the in depend-kind in that a task with a dependency with that depend-kind is dependent on all previously created sibling tasks with matching address unless they have the same depend-kind. In the in depend-kind case everything is dependent except for in -> in dependency, for inoutset everything is dependent except for inoutset -> inoutset dependency. mutexinoutset is also similar (everything is dependent except for mutexinoutset -> mutexinoutset dependency), but there is also the additional restriction that only one task with mutexinoutset for each address can be scheduled at once (i.e. mutual exclusitivty). For now we support mutexinoutset the same as inout/out, but the inoutset support is full. In order not to bump the ABI for dependencies each time (we've bumped it already once, the old ABI supports only inout/out and in depend-kind, the new ABI supports inout/out, mutexinoutset, in and depobj), this patch arranges for inoutset to be at least for the time being always handled as if it was specified through depobj even when it is not. So it uses the new ABI for that and inoutset are represented like depobj - pointer to a pair of pointers where the first one will be the actual address of the object mentioned in depend clause and second pointer will be (void *) GOMP_DEPEND_INOUTSET. 2022-05-17 Jakub Jelinek <jakub@redhat.com> gcc/ * tree-core.h (enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_INOUTSET. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_DEPEND_INOUTSET. * gimplify.cc (gimplify_omp_depend): Likewise. * omp-low.cc (lower_depend_clauses): Likewise. gcc/c-family/ * c-omp.cc (c_finish_omp_depobj): Handle OMP_CLAUSE_DEPEND_INOUTSET. gcc/c/ * c-parser.cc (c_parser_omp_clause_depend): Parse inoutset depend-kind. (c_parser_omp_depobj): Likewise. gcc/cp/ * parser.cc (cp_parser_omp_clause_depend): Parse inoutset depend-kind. (cp_parser_omp_depobj): Likewise. * cxx-pretty-print.cc (cxx_pretty_printer::statement): Handle OMP_CLAUSE_DEPEND_INOUTSET. gcc/testsuite/ * c-c++-common/gomp/all-memory-1.c (boo): Add test with inoutset depend-kind. * c-c++-common/gomp/all-memory-2.c (boo): Likewise. * c-c++-common/gomp/depobj-1.c (f1): Likewise. (f2): Adjusted expected diagnostics. * g++.dg/gomp/depobj-1.C (f4): Adjust expected diagnostics. include/ * gomp-constants.h (GOMP_DEPEND_INOUTSET): Define. libgomp/ * libgomp.h (struct gomp_task_depend_entry): Change is_in type from bool to unsigned char. * task.c (gomp_task_handle_depend): Handle GOMP_DEPEND_INOUTSET. Ignore dependencies where task->depend[i].is_in && task->depend[i].is_in == ent->is_in rather than just task->depend[i].is_in && ent->is_in. Remember whether GOMP_DEPEND_IN loop is needed and guard the loop with that conditional. (gomp_task_maybe_wait_for_dependencies): Handle GOMP_DEPEND_INOUTSET. Ignore dependencies where elem.is_in && elem.is_in == ent->is_in rather than just elem.is_in && ent->is_in. * testsuite/libgomp.c-c++-common/depend-1.c (test): Add task with inoutset depend-kind. * testsuite/libgomp.c-c++-common/depend-2.c (test): Likewise. * testsuite/libgomp.c-c++-common/depend-3.c (test): Likewise. * testsuite/libgomp.c-c++-common/depend-inoutset-1.c: New test. (cherry picked from commit 2c16eb3157f86ae561468c540caf8eb326106b5f)
2022-07-04OpenMP: Add omp_all_memory support to FortranTobias Burnus9-20/+231
Fortran part to the C/C++/backend implementation r13-337-g7f78783dbedca0183d193e475262ca3c489fd365 gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_namelist): Handle omp_all_memory. * openmp.cc (gfc_match_omp_variable_list, gfc_match_omp_depend_sink, gfc_match_omp_clauses, resolve_omp_clauses): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses, gfc_trans_omp_depobj): Likewise. * resolve.cc (resolve_symbol): Reject it as symbol. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1): Set omp_all_memory to 'Y'. * testsuite/libgomp.fortran/depend-5.f90: New test. * testsuite/libgomp.fortran/depend-6.f90: New test. * testsuite/libgomp.fortran/depend-7.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/all-memory-1.f90: New test. * gfortran.dg/gomp/all-memory-2.f90: New test. * gfortran.dg/gomp/all-memory-3.f90: New test. (cherry picked from commit 4f94c38a9237b728b3a3f76c169b5b47f6c45187)
2022-07-04OpenMP, C++: Add template support for the has_device_addr clause.Marcel Vollweiler3-2/+18
This patch adds support for list items in the has_device_addr clause which type is given by C++ template parameters. gcc/cp/ChangeLog: * pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR. * semantics.cc (finish_omp_clauses): Added template decl processing. libgomp/ChangeLog: * testsuite/libgomp.c++/target-has-device-addr-7.C: New test. * testsuite/libgomp.c++/target-has-device-addr-8.C: New test. * testsuite/libgomp.c++/target-has-device-addr-9.C: New test. (cherry picked from commit b4fb9f4f9a10d825302cfb5a0ecefa796570d8bc)
2022-07-04OpenMP/Fortran: Use firstprivat not alloc for ptr attach for arraysTobias Burnus2-3/+18
For a non-descriptor array, map(A(n:m)) was mapped as map(tofrom:A[n-1] [len: ...]) map(alloc:A [pointer assign, bias: ...]) with this patch, it is changed to map(tofrom:A[n-1] [len: ...]) map(firstprivate:A [pointer assign, bias: ...]) The latter avoids an alloc - and also avoids the race condition with nowait in the enclosed testcase. (Note: predantically, the testcase is invalid since OpenMP 5.1, violating the map clause restriction at [354:10-13]. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_trans_omp_clauses): When mapping nondescriptor array sections, use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER for the pointer attachment. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-nowait-array-section.f90: New test. (cherry picked from commit a46d6268371c446566f656858aada8775a0c988e)
2022-07-04Don't use gori depedencies to optimize.Andrew MacLeod2-13/+31
The routine fold_using_range::relation_fold_and_or needs to verify that both operands of 2 stmts are the same, and uses GORIs dependency cache for this. This cache cannot be counted on to reflect the current contents of a stmt, expecially in the presence of an IL changing pass. Instead, look at the statement operands. PR tree-optimization/106114 gcc/ * gimple-range-fold.cc (fold_using_range::relation_fold_and_or): Check statement operands instead of GORI cache. gcc/testsuite/ * gcc.dg/pr106114.c: New.
2022-07-04d: Fix error: aggregate value used where floating point was expectedIain Buclaw7-3/+178
Casting from vector to static array is permitted, and the frontend generates a reinterpret cast, but casting back the other way resulted in an error. This has been fixed to be properly handled in the code generation pass of VectorExp, and the conversion for lvalue and rvalue handling done in convert_expr and convert_for_rvalue respectively. PR d/106139 gcc/d/ChangeLog: * d-convert.cc (convert_expr): Handle casting from array to vector. (convert_for_rvalue): Rewrite vector to array casts of the same element type into a constructor. (convert_for_assignment): Return calling convert_for_rvalue. * expr.cc (ExprVisitor::visit (VectorExp *)): Handle generating a vector expression from a static array. * toir.cc (IRVisitor::visit (ReturnStatement *)): Call convert_for_rvalue on return value. gcc/testsuite/ChangeLog: * gdc.dg/pr106139a.d: New test. * gdc.dg/pr106139b.d: New test. * gdc.dg/pr106139c.d: New test. * gdc.dg/pr106139d.d: New test. (cherry picked from commit 329bef49da30158d30fed1106002bb71674776bd)
2022-07-04PR target/105991: Recognize PLUS and XOR forms of rldimi in rs6000.md.Roger Sayle2-1/+32
This patch addresses PR target/105991 where a change to prefer representing shifts and adds at the tree-level as multiplications, causes problems for the rldimi patterns in the powerpc backend. The issue is that rs6000.md models this pattern using IOR, and some variants that have the equivalent PLUS or XOR in the RTL fail to match some *rotl<mode>4_insert patterns. This is fixed in this patch by adding a define_insn_and_split to locally canonicalize the PLUS and XOR forms to the backend's preferred IOR form. Backported from master. 2022-07-04 Roger Sayle <roger@nextmovesoftware.com> Marek Polacek <polacek@redhat.com> Segher Boessenkool <segher@kernel.crashing.org> Kewen Lin <linkw@linux.ibm.com> gcc/ChangeLog PR target/105991 * config/rs6000/rs6000.md (rotl<mode>3_insert_3): Check that exact_log2 doesn't return -1 (or zero). (plus_xor): New code iterator. (*rotl<mode>3_insert_3_<code>): New define_insn_and_split. gcc/testsuite/ChangeLog PR target/105991 * gcc.target/powerpc/pr105991.c: New test case.
2022-07-04loongarch: use -mno-check-zero-division as the default for optimized codeXi Ruoyao4-7/+18
Integer division by zero is undefined behavior anyway, and there are already many platforms where neither the GCC port and the hardware do anything to trap on division by zero. So any portable program shall not rely on SIGFPE on division by zero, in both theory and practice. As the result, there is no real reason to cost two additional instructions just for the trap on division by zero with a new ISA. One remaining reason to trap on division by zero may be debugging, especially while -fsanitize=integer-divide-by-zero is not implemented for LoongArch yet. To make debugging easier, keep -mcheck-zero-division as the default for -O0 and -Og, but use -mno-check-zero-division as the default for all other optimization levels. Backport this behavior change for 12.2, so 12.1 will be the only release with a different default. Co-authored-by: Lulu Cheng <chenglulu@loongson.cn> gcc/ChangeLog: * config/loongarch/loongarch.cc (loongarch_check_zero_div_p): New static function. (loongarch_idiv_insns): Use loongarch_check_zero_div_p instead of TARGET_CHECK_ZERO_DIV. (loongarch_output_division): Likewise. * common/config/loongarch/loongarch-common.cc (TARGET_DEFAULT_TARGET_FLAGS): Remove unneeded hook. * doc/invoke.texi: Update to match the new behavior. gcc/testsuite/ChangeLog: * gcc.c-torture/execute/20101011-1.c (dg-additional-options): add -mcheck-zero-division for LoongArch targets. (cherry picked from commit f150dc1bd11802b70277f0fa209f2d23695a1095)
2022-07-04Daily bump.GCC Administrator4-1/+58
2022-07-03Fortran: error recovery on invalid CLASS(), PARAMETER declarations [PR105243]Harald Anlauf3-1/+26
gcc/fortran/ChangeLog: PR fortran/103137 PR fortran/103138 PR fortran/103693 PR fortran/105243 * decl.cc (gfc_match_data_decl): Reject CLASS entity declaration when it is given the PARAMETER attribute. gcc/testsuite/ChangeLog: PR fortran/103137 PR fortran/103138 PR fortran/103693 PR fortran/105243 * gfortran.dg/class_58.f90: Fix test. * gfortran.dg/class_73.f90: New test. Co-authored-by: Steven G. Kargl <kargl@gcc.gnu.org> (cherry picked from commit 4c233cabbe388a6b8957c1507e129090e9267ceb)
2022-07-03Fortran: improve error recovery for EXTENDS_TYPE_OF() [PR106121]Harald Anlauf2-0/+24
gcc/fortran/ChangeLog: PR fortran/106121 * simplify.cc (gfc_simplify_extends_type_of): Do not attempt to simplify when one of the arguments is a CLASS variable that was not properly declared. gcc/testsuite/ChangeLog: PR fortran/106121 * gfortran.dg/extends_type_of_4.f90: New test. Co-authored-by: Steven G. Kargl <kargl@gcc.gnu.org> (cherry picked from commit b8f284d3673004dffae714b56ed663467c2a52a7)
2022-07-03PR target/106122: Don't update %esp via the stack with -Oz on x86.Roger Sayle2-1/+17
When optimizing for size with -Oz, setting a register can be minimized by pushing an immediate value to the stack and popping it to the destination. Alas the one general register that shouldn't be updated via the stack is the stack pointer itself, where "pop %esp" can't be represented in GCC's RTL ("use of a register mentioned in pre_inc, pre_dec, post_inc or post_dec is not permitted within the same instruction"). This patch fixes PR target/106122 by explicitly checking for SP_REG in the problematic peephole2. 2022-07-O3 Roger Sayle <roger@nextmovesoftware.com> gcc/ChangeLog PR target/106122 * config/i386/i386.md (peephole2): Avoid generating pop %esp when optimizing for size. gcc/testsuite/ChangeLog PR target/106122 * gcc.target/i386/pr106122.c: New test case.
2022-07-03Daily bump.GCC Administrator4-1/+30
2022-07-02jit: avoid calloc() poisoning on musl [PR106102]Sergei Trofimovich4-4/+7
On musl <pthread.h> uses calloc() (via <sched.h>). jit/ includes it directly and exposes use of poisoned calloc(): /build/build/./prev-gcc/xg++ ... ../../gcc-13-20220626/gcc/jit/jit-playback.cc make[3]: *** [Makefile:1143: jit/libgccjit.o] Error 1 make[3]: *** Waiting for unfinished jobs.... In file included from /<<NIX>>/musl-1.2.3-dev/include/pthread.h:30, from ../../gcc-13-20220626/gcc/jit/jit-playback.cc:44: /<<NIX>>/musl-1.2.3-dev/include/sched.h:84:7: error: attempt to use poisoned "calloc" 84 | void *calloc(size_t, size_t); | ^ /<<NIX>>/musl-1.2.3-dev/include/sched.h:124:36: error: attempt to use poisoned "calloc" 124 | #define CPU_ALLOC(n) ((cpu_set_t *)calloc(1,CPU_ALLOC_SIZE(n))) | ^ The change moves <pthread.h> inclusion to "system.h" under new INCLUDE_PTHREAD_H guard and uses this mechanism in libgccjit. gcc/ PR c++/106102 * system.h: Introduce INCLUDE_PTHREAD_H macros to include <pthread.h>. gcc/jit/ PR c++/106102 * jit-playback.cc: Include <pthread.h> via "system.h" to avoid calloc() poisoning. * jit-recording.cc: Ditto. * libgccjit.cc: Ditto. (cherry picked from commit 49d508065bdd36fb1a9b6aad9666b1edb5e06474)
2022-07-02c++: avoid <memory> poisoning on musl [PR106102]Sergei Trofimovich3-0/+3
On musl <pthread.h> uses calloc() (via <sched.h>). <memory> includes it indirectly and exposes use of poisoned calloc() when module code is built: /build/build/./prev-gcc/xg++ ... ../../gcc-13-20220626/gcc/cp/mapper-resolver.cc In file included from /<<NIX>>/musl-1.2.3-dev/include/pthread.h:30, from /build/build/prev-x86_64-unknown-linux-musl/libstdc++-v3/include/x86_64-unknown-linux-musl/bits/gthr-default.h:35, .... from /build/build/prev-x86_64-unknown-linux-musl/libstdc++-v3/include/memory:77, from ../../gcc-13-20220626/gcc/../libcody/cody.hh:24, from ../../gcc-13-20220626/gcc/cp/../../c++tools/resolver.h:25, from ../../gcc-13-20220626/gcc/cp/../../c++tools/resolver.cc:23, from ../../gcc-13-20220626/gcc/cp/mapper-resolver.cc:32: /<<NIX>>/musl-1.2.3-dev/include/sched.h:84:7: error: attempt to use poisoned "calloc" 84 | void *calloc(size_t, size_t); | ^ /<<NIX>>/musl-1.2.3-dev/include/sched.h:124:36: error: attempt to use poisoned "calloc" 124 | #define CPU_ALLOC(n) ((cpu_set_t *)calloc(1,CPU_ALLOC_SIZE(n))) | ^ gcc/cp/ PR c++/106102 * mapper-client.cc: Include <memory> via "system.h". * mapper-resolver.cc: Ditto. * module.cc: Ditto. libcc1/ PR c++/106102 * libcc1plugin.cc: Include <memory> via "system.h". * libcp1plugin.cc: Ditto. (cherry picked from commit 3b21c21f3f5726823e19728fdd1571a14aae0fb3)
2022-07-02Daily bump.GCC Administrator4-1/+45
2022-07-01Merge branch 'releases/gcc-12' into devel/omp/gcc-12Kwok Cheung Yeung296-85270/+89454
Merged up to commit e748398b3ef6412ef35b85ef6b0893809aeb49cd (1st July 2022).
2022-07-01amdgcn, openmp: Auto-detect USM mode and set HSA_XNACKAndrew Stubbs4-2/+78
The AMD GCN runtime must be set to the correct mode for Unified Shared Memory to work, but this is not always clear at compile and link time due to the split nature of the offload compilation pipeline. This patch sets a new attribute on OpenMP offload functions to ensure that the information is passed all the way to the backend. The backend then places a marker in the assembler code for mkoffload to find. Finally mkoffload places a constructor function into the final program to ensure that the HSA_XNACK environment variable passes the correct mode to the GPU. The HSA_XNACK variable must be set before the HSA runtime is even loaded, so it makes more sense to have this set within the constructor than at some point later within libgomp or the GCN plugin. gcc/ChangeLog: * config/gcn/gcn.c (unified_shared_memory_enabled): New variable. (gcn_init_cumulative_args): Handle attribute "omp unified memory". (gcn_hsa_declare_function_name): Emit "MKOFFLOAD OPTIONS: USM+". * config/gcn/mkoffload.c (TEST_XNACK_OFF): New macro. (process_asm): Detect "MKOFFLOAD OPTIONS: USM+". Emit configure_xnack constructor, as required. * omp-low.c (create_omp_child_function): Add attribute "omp unified memory".
2022-07-01amdgcn: Support XNACK modeAndrew Stubbs8-86/+171
The XNACK feature allows memory load instructions to restart safely following a page-miss interrupt. This is useful for shared-memory devices, like APUs, and to implement OpenMP Unified Shared Memory. To support the feature we must be able to set the appropriate meta-data and set the load instructions to early-clobber. When the port supports scheduling of s_waitcnt instructions there will be further requirements. gcc/ChangeLog: * config/gcn/gcn-hsa.h (XNACKOPT): New macro. (ASM_SPEC): Use XNACKOPT. * config/gcn/gcn-opts.h (enum sram_ecc_type): Rename to ... (enum hsaco_attr_type): ... this, and generalize the names. (TARGET_XNACK): New macro. * config/gcn/gcn-valu.md (gather<mode>_insn_1offset<exec>): Add xnack compatible alternatives. (gather<mode>_insn_2offsets<exec>): Likewise. * config/gcn/gcn.c (gcn_option_override): Permit -mxnack for devices other than Fiji. (gcn_expand_epilogue): Remove early-clobber problems. (output_file_start): Emit xnack attributes. (gcn_hsa_declare_function_name): Obey -mxnack setting. * config/gcn/gcn.md (xnack): New attribute. (enabled): Rework to include "xnack" attribute. (*movbi): Add xnack compatible alternatives. (*mov<mode>_insn): Likewise. (*mov<mode>_insn): Likewise. (*mov<mode>_insn): Likewise. (*movti_insn): Likewise. * config/gcn/gcn.opt (-mxnack): Add the "on/off/any" syntax. (sram_ecc_type): Rename to ... (hsaco_attr_type: ... this.) * config/gcn/mkoffload.c (SET_XNACK_ANY): New macro. (TEST_XNACK): Delete. (TEST_XNACK_ANY): New macro. (TEST_XNACK_ON): New macro. (main): Support the new -mxnack=on/off/any syntax.
2022-07-01c++: simpler fix for PR106024Jason Merrill1-20/+2
Actually, for release branches let's just avoid the lookup for the lambdas that are the problematic case and only make the bigger change on trunk. PR c++/106024 gcc/cp/ChangeLog: * parser.cc (cp_parser_lookup_name): Limit previous change to lambdas.
2022-07-01Fix ICE on shVladimir Makarov1-0/+6
gcc/ PR target/103722 * config/sh/sh.cc (sh_register_move_cost): Avoid cost "2" (which is special) for various scenarios. (cherry picked from commit ce1580252ea57de23a595e9804ea87ed4353aa6a)
2022-07-01c++: lambda template in requires [PR105541]Jason Merrill3-6/+15
Since the patch for PR103408, the template parameters for the lambda in this test have level 1 instead of 2, and we were treating null template args as 1 level of arguments, so tsubst_template_parms decided it had nothing to do. Fixed by distinguishing between <> and no args at all, which is what we have in our "substitution" in a requires-expression. PR c++/105541 gcc/cp/ChangeLog: * cp-tree.h (TMPL_ARGS_DEPTH): 0 for null args. * parser.cc (cp_parser_enclosed_template_argument_list): Use 0-length TREE_VEC for <>. gcc/testsuite/ChangeLog: * g++.dg/cpp2a/lambda-requires1.C: New test.
2022-07-01c++: tweak resolve_args changeJason Merrill1-1/+1
I don't know why I used tf_error instead of complain here. PR c++/105779 gcc/cp/ChangeLog: * call.cc (resolve_args): Use complain.
2022-07-01c++: dependent generic lambda template-id [PR106024]Jason Merrill2-1/+36
We were wrongly looking up the generic lambda op() in a dependent scope, and then trying to look up its instantiation at substitution time, but lambdas aren't instantiated, so we crashed. The fix is to not look into dependent class scopes. But this created trouble with wrongly trying to use a template from the enclosing scope when we aren't actually looking at a template-argument-list, in template/lookup18.C, so let's avoid that. PR c++/106024 gcc/cp/ChangeLog: * parser.cc (missing_template_diag): Factor out... (cp_parser_id_expression): ...from here. (cp_parser_lookup_name): Don't look in dependent object_type. gcc/testsuite/ChangeLog: * g++.dg/cpp2a/lambda-generic10.C: New test.
2022-07-01Fix gfortran.dg/gomp/num-teams-2.f90Tobias Burnus2-6/+10
OG11 contrary to mainline issues an error for resolve_positive_int_expr (-> OG11 commit a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c). Update testcase accordingly. gcc/testsuite/ * gfortran.dg/gomp/num-teams-2.f90: Use dg-error not dg-warning.
2022-07-01Daily bump.GCC Administrator3-1/+55
2022-06-30Fortran: handle explicit-shape specs with constant bounds [PR105954]Harald Anlauf2-0/+38
gcc/fortran/ChangeLog: PR fortran/105954 * decl.cc (variable_decl): Adjust upper bounds for explicit-shape specs with constant bound expressions to ensure non-negative extents. gcc/testsuite/ChangeLog: PR fortran/105954 * gfortran.dg/pr105954.f90: New test. (cherry picked from commit a312407bd715647f7c11b67e0a52effc94d0f15d)
generated by cgit v1.1 at 2025-09-23 09:43:09 +0000