mirror of
https://forge.sourceware.org/marek/gcc.git
synced 2026-02-22 03:47:02 -05:00
reflection
1360 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
Andrew Stubbs
|
57af002207 |
amdgcn: remove gfx803 "Fiji" support
The gfx803 "Fiji" device was deprecated in GCC 14, removed from LLVM 18, and
hasn't worked properly with the drivers since about ROCm 4.
This patch removes the device from GCC options and documentation, and removes
the direct mentions from the internals.
The TARGET_GCN3 support in the back-end is now unused and can be removed (in a
follow-up patch).
gcc/ChangeLog:
* config.gcc (amdgcn-*-*): Remove "fiji" from with_arch checks.
* config/gcn/gcn-hsa.h (ABI_VERSION_SPEC): Remove fiji alternative.
(NO_XNACK): Likewise.
(NO_SRAM_ECC): Likewise.
(ASM_SPEC): Remove "%{}" around ABI_VERSION_SPEC.
* config/gcn/gcn-opts.h (enum processor_type): Remove PROCESSOR_FIJI.
(TARGET_FIJI): Delete.
* config/gcn/gcn.cc (gcn_option_override): Remove Fiji.
(gcn_omp_device_kind_arch_isa): Likewise.
(output_file_start): Likewise.
* config/gcn/gcn.h (TARGET_CPU_CPP_BUILTINS): Likewise.
* config/gcn/gcn.opt (gpu_type): Likewise.
(march, mtune): Change default to PROCESSOR_VEGA10.
* config/gcn/mkoffload.cc (EF_AMDGPU_MACH_AMDGCN_GFX803): Delete.
(copy_early_debug_info): Remove elf_flags_actual.
Use ELFABIVERSION_AMDGPU_HSA_V4 unconditionally.
(get_arch): Remove Fiji.
(main): Remove gfx803.
* config/gcn/t-omp-device
(omp-device-properties-gcn): Remove fiji and gfx803.
* doc/install.texi (amdgcn*-*-*): Remove fiji and special instructions.
* doc/invoke.texi: Remove fiji.
libgomp/ChangeLog:
* libgomp.texi: Remove fiji and gfx803.
* testsuite/libgomp.c/declare-variant-4.h: Remove fiji and gfx803.
* testsuite/libgomp.c/declare-variant-4-fiji.c: Removed.
* testsuite/libgomp.c/declare-variant-4-gfx803.c: Removed.
|
||
|
Alex Coplan
|
e4d3e7f9ad |
testsuite: Rename scanltranstree.exp -> scanltrans.exp
Since r15-3254-g3f51f0dc88ec21c1ec79df694200f10ef85915f4 added scan-ltrans-rtl* variants to scanltranstree.exp, it no longer makes sense to have "tree" in the name. This renames the file accordingly and updates users. libatomic/ChangeLog: * testsuite/lib/libatomic.exp: Load scanltrans.exp instead of scanltranstree.exp. libgomp/ChangeLog: * testsuite/lib/libgomp.exp: Load scanltrans.exp instead of scanltranstree.exp. libitm/ChangeLog: * testsuite/lib/libitm.exp: Load scanltrans.exp instead of scanltranstree.exp. libphobos/ChangeLog: * testsuite/lib/libphobos-dg.exp: Load scanltrans.exp instead of scanltranstree.exp. libvtv/ChangeLog: * testsuite/lib/libvtv.exp: Load scanltrans.exp instead of scanltranstree.exp. gcc/testsuite/ChangeLog: * gcc.dg-selftests/dg-final.exp: Load scanltrans.exp instead of scanltranstree.exp. * lib/gcc-dg.exp: Likewise. * lib/scanltranstree.exp: Rename to ... * lib/scanltrans.exp: ... this. |
||
|
Tobias Burnus
|
0beac1db38 |
libgomp: Add interop types and routines to OpenMP's headers and module
This commit adds OpenMP 5.1+'s interop enumeration, type and routine declarations to the C/C++ header file and, new in OpenMP TR13, also to the Fortran module and omp_lib.h header file. While a stub implementation is provided, only with foreign runtime support by the libgomp GPU plugins and with the 'interop' directive, this becomes really useful. libgomp/ChangeLog: * fortran.c (omp_get_interop_str_, omp_get_interop_name_, omp_get_interop_type_desc_, omp_get_interop_rc_desc_): Add. * libgomp.map (GOMP_5.1.3): New; add interop routines. * omp.h.in: Add interop typedefs, enum and prototypes. (__GOMP_DEFAULT_NULL): Define. (omp_target_memcpy_async, omp_target_memcpy_rect_async): Use it for the optional depend argument. * omp_lib.f90.in: Add paramters and interfaces for interop. * omp_lib.h.in: Likewise; move F90 '&' to column 81 for -ffree-length-80. * target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/gcn/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/nvptx/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * testsuite/libgomp.c-c++-common/interop-routines-1.c: New test. * testsuite/libgomp.c-c++-common/interop-routines-2.c: New test. * testsuite/libgomp.fortran/interop-routines-1.F90: New test. * testsuite/libgomp.fortran/interop-routines-2.F90: New test. * testsuite/libgomp.fortran/interop-routines-3.F: New test. * testsuite/libgomp.fortran/interop-routines-4.F: New test. * testsuite/libgomp.fortran/interop-routines-5.F: New test. * testsuite/libgomp.fortran/interop-routines-6.F: New test. * testsuite/libgomp.fortran/interop-routines-7.F90: New test. |
||
|
Thomas Schwinge
|
9f5d22e3e2 |
OpenMP: Constructors and destructors for "declare target" static aggregates: Fix effective-target keyword in test cases
(Most of) the tests added in commit
|
||
|
Tobias Burnus
|
e3a6dec326 |
libgomp.c++/static-aggr-constructor-destructor-{1,2}.C: Fix scan-tree-dump
In principle, the optimized dump should be the same on the host, but as 'nohost' is not handled, is is present. However when ENABLE_OFFLOADING is false, it is handled early enough to remove the function. libgomp/ChangeLog: * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: Split scan-tree-dump into with and without target offload_target_any. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: Likewise. |
||
|
Tobias Burnus
|
f1bfba3a9b |
OpenMP: Constructors and destructors for "declare target" static aggregates
This commit also compile-time expands (__builtin_)omp_is_initial_device for both Fortran and C/C++ (unless, -fno-builtin-omp_is_initial_device is used). But the main change is: This commit adds support for running constructors and destructors for static (file-scope) aggregates for C++ objects which are marked with "declare target" directives on OpenMP offload targets. Before this commit, space is allocated on the target for such aggregates, but nothing ever constructs them properly, so they end up zero-initialised. (See the new test static-aggr-constructor-destructor-3.C for a reason why running constructors on the target is preferable to e.g. constructing on the host and then copying the resulting object to the target.) 2024-08-07 Julian Brown <julian@codesourcery.com> Tobias Burnus <tobias@baylibre.com> gcc/ChangeLog: * builtins.def (DEF_GOMP_BUILTIN_COMPILER): Define DEF_GOMP_BUILTIN_COMPILER to handle the non-prefix version. * gimple-fold.cc (gimple_fold_builtin_omp_is_initial_device): New. (gimple_fold_builtin): Call it. * omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): Define. * tree.cc (get_file_function_name): Support names for on-target constructor/destructor functions. gcc/cp/ * decl2.cc (tree-inline.h): Include. (static_init_fini_fns): Bump to four entries. Update comment. (start_objects, start_partial_init_fini_fn): Add 'omp_target' parameter. Support "declare target" decls. Update forward declaration. (emit_partial_init_fini_fn): Add 'host_fn' parameter. Return tree for the created function. Support "declare target". (OMP_SSDF_IDENTIFIER): New macro. (partition_vars_for_init_fini): Support partitioning "declare target" variables also. (generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support "declare target" decls. (c_parse_final_cleanups): Support constructors/destructors on OpenMP offload targets. gcc/fortran/ChangeLog: * gfortran.h (gfc_option_t): Add disable_omp_is_initial_device. * lang.opt (fbuiltin-): Add. * options.cc (gfc_handle_option): Handle -fno-builtin-omp_is_initial_device. * f95-lang.cc (gfc_init_builtin_functions): Handle DEF_GOMP_BUILTIN_COMPILER. * trans-decl.cc (gfc_get_extern_function_decl): Add code to use DEF_GOMP_BUILTIN_COMPILER for 'omp_is_initial_device'. libgomp/ChangeLog: * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host.c: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: New test. * testsuite/libgomp.fortran/target-is-initial-host.f: New test. * testsuite/libgomp.fortran/target-is-initial-host.f90: New test. * testsuite/libgomp.fortran/target-is-initial-host-2.f90: New test. Co-authored-by: Tobias Burnus <tobias@baylibre.com> |
||
|
Tobias Burnus
|
aa689684d2 |
libgomp.c-c++-common/target-link-2.c: Fix test on multi-device systems
libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/target-link-2.c: Reset variable value to handle multi-device tests. |
||
|
Tobias Burnus
|
c99cdcab4f |
omp-offload.cc: Fix value-expr handling of 'declare target link' vars [PR115637]
As the PR and included testcase shows, replacing 'arr2' by its value expression
'*arr2$13$linkptr' failed for
MEM <uint128_t> [(c_char * {ref-all})&arr2]
which left 'arr2' in the code as unknown symbol. Now expand the value expression
already in pass_omp_target_link::execute's process_link_var_op walk_gimple_stmt
walk - and don't rely on gimple_regimplify_operands.
PR middle-end/115637
gcc/ChangeLog:
* gimplify.cc (gimplify_body): Fix macro name in the comment.
* omp-offload.cc (find_link_var_op): Rename to ...
(process_link_var_op): ... this. Replace value expr.
(pass_omp_target_link::execute): Update walk_gimple_stmt call.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/declare-target-link.f90: Uncomment
now working code.
Co-authored-by: Richard Biener <rguenther@suse.de
|
||
|
Sam James
|
9ad3d1c581 |
testsuite: fix dg-require-effective-target order vs dg-additional-sources
Per gccint, 'dg-require-effective-target' must come before any 'dg-additional-sources' directives. Fix a handful of deviant cases. gcc/testsuite/ChangeLog: * gcc.target/aarch64/aapcs64/func-ret-3.c: Fix dg-require-effective-target directive order. * gcc.target/aarch64/aapcs64/func-ret-4.c: Likewise. * gfortran.dg/PR100914.f90: Likewise. libgomp/ChangeLog: * testsuite/libgomp.c++/pr24455.C: Fix dg-require-effective-target directive order. * testsuite/libgomp.c/pr24455.c: Likewise. |
||
|
Sam James
|
1c85b16f19 |
testsuite: libgomp: fix dg-do run typo
'dg-run' is not a valid dejagnu directive, 'dg-do run' is needed here for the test to be executed. That said, it actually seems to be executed for me anyway, presumably a default in the directory, but let's fix it to be consistent with other uses in the tree and in that test directory even. libgomp/ChangeLog: * testsuite/libgomp.c++/declare-target-indirect-1.C: Fix 'dg-run' typo. |
||
|
Tobias Burnus
|
29b1587e7d |
OpenMP/Fortran: Fix handling of 'declare target' with 'link' clause [PR115559]
Contrary to a normal 'declare target', the 'declare target link' attribute also needs to set node->offloadable and push the offload_vars in the front end. Linked variables require that the data is mapped. For module variables, this can happen anywhere. For variables in an external subprograms or the main programm, this can only happen in the either that program itself or in an internal subprogram. - Whether a variable is just normally mapped or linked then becomes relevant if a device routine exists that can access that variable, i.e. an internal procedure has then to be marked as declare target. PR fortran/115559 gcc/fortran/ChangeLog: * trans-common.cc (build_common_decl): Add 'omp declare target' and 'omp declare target link' variables to offload_vars. * trans-decl.cc (add_attributes_to_decl): Likewise; update args and call decl_attributes. (get_proc_pointer_decl, gfc_get_extern_function_decl, build_function_decl): Update calls. (gfc_get_symbol_decl): Likewise; move after 'DECL_STATIC (t)=1' to avoid errors with symtab_node::get_create. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-link.f90: New test. |
||
|
Tobias Burnus
|
14c47e7eb0 |
libgomp: Fix declare target link with offset array-section mapping [PR116107]
Assume that 'int var[100]' is 'omp declare target link(var)'. When now mapping an array section with offset such as 'map(to:var[20:10])', the device-side link pointer has to store &<device-storage-data>[0] minus the offset such that var[20] will access <device-storage-data>[0]. But the offset calculation was missed such that the device-side 'var' pointed to the first element of the mapped data - and var[20] points beyond at some invalid memory. PR middle-end/116107 libgomp/ChangeLog: * target.c (gomp_map_vars_internal): Honor array mapping offsets with declare-target 'link' variables. * testsuite/libgomp.c-c++-common/target-link-2.c: New test. |
||
|
Paul Thomas
|
8d6994f33a |
libgomp: Remove bogus warnings from privatized-ref-2.f90.
2024-07-19 Paul Thomas <pault@gcc.gnu.org> libgomp/ChangeLog * testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Cut dg-note about 'a' and remove bogus warnings about its array descriptor components being used uninitialized. |
||
|
Andrew Stubbs
|
64001441ec |
libgomp, openmp: Add ompx_gnu_pinned_mem_alloc
This creates a new predefined allocator as a shortcut for using pinned memory with OpenMP. This is not in the OpenMP standard so it uses the "ompx" namespace and an independent enum baseline of 200 (selected to not clash with other known implementations). The allocator is equivalent to using a custom allocator with the pinned trait and the null fallback trait. One motivation for having this feature is for use by the (planned) -foffload-memory=pinned feature. gcc/fortran/ChangeLog: * openmp.cc (is_predefined_allocator): Update valid ranges to incorporate ompx_gnu_pinned_mem_alloc. libgomp/ChangeLog: * allocator.c (ompx_gnu_min_predefined_alloc): New. (ompx_gnu_max_predefined_alloc): New. (predefined_alloc_mapping): Rename to ... (predefined_omp_alloc_mapping): ... this. (predefined_ompx_gnu_alloc_mapping): New. (_Static_assert): Adjust for the new name, and add a new assert for the new table. (predefined_allocator_p): New. (predefined_alloc_mapping): New. (omp_aligned_alloc): Support ompx_gnu_pinned_mem_alloc. Use predefined_allocator_p and predefined_alloc_mapping. (omp_free): Likewise. (omp_alligned_calloc): Likewise. (omp_realloc): Likewise. * env.c (parse_allocator): Add ompx_gnu_pinned_mem_alloc. * libgomp.texi: Document ompx_gnu_pinned_mem_alloc. * omp.h.in (omp_allocator_handle_t): Add ompx_gnu_pinned_mem_alloc. * omp_lib.f90.in: Add ompx_gnu_pinned_mem_alloc. * omp_lib.h.in: Add ompx_gnu_pinned_mem_alloc. * testsuite/libgomp.c/alloc-pinned-5.c: New test. * testsuite/libgomp.c/alloc-pinned-6.c: New test. * testsuite/libgomp.fortran/alloc-pinned-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/allocate-pinned-1.f90: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> |
||
|
Andrew Stubbs
|
90efaebf95 |
libgomp: change alloc-pinned tests failure mode
The feature doesn't work on non-Linux hosts, at present, so skip the tests entirely. On Linux systems that have insufficient lockable memory configured we still need to fail or else the feature won't be getting tested when we think it is, but now there's a message to explain why. libgomp/ChangeLog: * testsuite/libgomp.c/alloc-pinned-1.c: Change dg-xfail-run-if to dg-skip-if. Correct spelling mistake. Abort on insufficient lockable memory. Use #error on non-linux hosts. * testsuite/libgomp.c/alloc-pinned-2.c: Likewise. |
||
|
Thomas Schwinge
|
3a4775d440 |
nvptx, libgfortran: Switch out of "minimal" mode
..., in order to enable (portions of) Fortran I/O, for example. libgfortran/ * configure.ac: No longer set 'LIBGFOR_MINIMAL' for nvptx. * configure: Regenerate. libgomp/ * libgomp.texi (nvptx): Update. * testsuite/libgomp.fortran/target-print-1-nvptx.f90: Remove. * testsuite/libgomp.fortran/target-print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/error_stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Adjust. * testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Adjust. * testsuite/libgomp.oacc-fortran/print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/stop-2.f: Adjust. Co-authored-by: Andrew Stubbs <ams@gcc.gnu.org> |
||
|
Thomas Schwinge
|
395ac0417a |
Clean up after newlib "nvptx: In offloading execution, map '_exit' to 'abort' [GCC PR85463]"
PR target/85463 libgfortran/ * runtime/minimal.c [__nvptx__] (exit): Don't override. libgomp/ * config/nvptx/error.c (exit): Don't override. * testsuite/libgomp.oacc-fortran/error_stop-1.f: Update. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise. * testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-1.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-2.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-3.f: Likewise. |
||
|
Jakub Jelinek
|
804c0f35a6 |
openmp: OpenMP loop transformation support
This patch is largely rewritten version of the https://gcc.gnu.org/pipermail/gcc-patches/2023-October/631764.html patch set which I've promissed to adjust the way I'd like it but didn't get to it until now. The previous series together in diffstat was 176 files changed, 12107 insertions(+), 298 deletions(-) This patch is 197 files changed, 10843 insertions(+), 212 deletions(-) and diff between the old series and new patch is 268 files changed, 8053 insertions(+), 9231 deletions(-) Only the 5.1/5.2 tile/unroll constructs are supported, in various places some preparations for the other 6.0 loop transformations constructs (interchange/reverse/fuse) are done, but certainly not complete and not everywhere. The important difference is that because tile/unroll partial map 1:1 the original loops to generated canonical loops and add another set of generated loops without canonical form inside of it, the tile/unroll partial constructs are terminal for the generated loop, one can't have some loops from the tile or unroll partial and some further loops from inside the body of that construct. The GENERIC representation attempts to match what the standard specifies, so there are separate OMP_TILE and OMP_UNROLL trees. If for a particular loop in a loop nest of some OpenMP loop it awaits a generated loop from a nested loop, or if in OMP_LOOPXFORM_LOWERED OMP_TILE/UNROLL construct a generated loop has been moved to some surrounding construct, that particular loop is represented by all NULL_TREEs in the OMP_FOR_{INIT,COND,INCR,ORIG_DECLS} vector. The lowering of the loop transforming constructs is done at gimplification time, at the start of gimplify_omp_for. I think this way it is more maintainable over magic clauses with various loop depths on the other looping constructs or the magic OMP_LOOP_TRANS construct. Though, I admit I'm still undecided how to represent the OpenMP 6.0 loop transformation case of say: #pragma omp for collapse (4) for (int i = 0; i < 32; ++i) #pragma omp interchange permutation (2, 1) #pragma omp reverse for (int j = 0; j < 32; ++j) #pragma omp reverse for (int k = 0; k < 32; ++k) for (int l = 0; l < 32; ++l) ; Surely the i loop would go to first vector elements of OMP_FOR_* of the work-sharing loop, then 2 loops are expecting generated loops from interchange which would be inside of the body. But the innermost l loop isn't part of the interchange, so the question is where to put it. One possibility is to have it in the 4th loop of the OMP_FOR, another possibility would be to add some artificial construct inside of the OMP_INTERCHANGE and 2 OMP_REVERSE bodies which would contain the inner loop(s), e.g. it could be OMP_INTERCHANGE without permutation clause or some artificial ones or whatever. I've recently raised various unclear things in the 5.1/5.2/TRs versions regarding loop transformations, in particular https://github.com/OpenMP/spec/issues/3908 https://github.com/OpenMP/spec/issues/3909 (sorry, private links unless you have OpenMP membership). Until those are resolved, I have a sorry on trying to mix generated loops with non-rectangular loops (way too many questions need to be answered before that can be done) and similarly for mixing non-perfectly nested loops with generated loops (again, it can be implemented somehow, but is way too unclear). The second issue is mostly about data sharing, which is ambiguous, the patch makes the artificial iterators of the loops effectively private in the associated constructs (more like local), but for user iterators doesn't do anything in particular, so for now one needs to use explicit data sharing clauses on the non-loop transformation OpenMP looping constructs or surrounding parallel/task/target etc. 2024-06-05 Jakub Jelinek <jakub@redhat.com> Frederik Harwath <frederik@codesourcery.com> Sandra Loosemore <sandra@codesourcery.com> gcc/ * tree.def (OMP_TILE, OMP_UNROLL): New tree codes. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_PARTIAL, OMP_CLAUSE_FULL and OMP_CLAUSE_SIZES. * tree.h (OMP_LOOPXFORM_CHECK): Define. (OMP_LOOPXFORM_LOWERED): Define. (OMP_CLAUSE_PARTIAL_EXPR): Define. (OMP_CLAUSE_SIZES_LIST): Define. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add entries for OMP_CLAUSE_{PARTIAL,FULL,SIZES}. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_{PARTIAL,FULL,SIZES}. (dump_generic_node): Handle OMP_TILE and OMP_UNROLL. Skip printing loops with NULL OMP_FOR_INIT (node) vector element. * gimplify.cc (is_gimple_stmt): Handle OMP_TILE and OMP_UNROLL. (gimplify_omp_taskloop_expr): For SAVE_EXPR use gimplify_save_expr. (gimplify_omp_loop_xform): New function. (gimplify_omp_for): Call omp_maybe_apply_loop_xforms and if that reshuffles what the passed pointer points to, retry or return GS_OK. Handle OMP_TILE and OMP_UNROLL. (gimplify_omp_loop): Call omp_maybe_apply_loop_xforms and if that reshuffles what the passed pointer points to, return GS_OK. (gimplify_expr): Handle OMP_TILE and OMP_UNROLL. * omp-general.h (omp_loop_number_of_iterations, omp_maybe_apply_loop_xforms): Declare. * omp-general.cc (omp_adjust_for_condition): For LE_EXPR and GE_EXPR with pointers, don't add/subtract one, but the size of what the pointer points to. (omp_loop_number_of_iterations, omp_apply_tile, find_nested_loop_xform, omp_maybe_apply_loop_xforms): New functions. gcc/c-family/ * c-common.h (c_omp_find_generated_loop): Declare. * c-gimplify.cc (c_genericize_control_stmt): Handle OMP_TILE and OMP_UNROLL. * c-omp.cc (c_finish_omp_for): Handle generated loops. (c_omp_is_loop_iterator): Likewise. (c_find_nested_loop_xform_r, c_omp_find_generated_loop): New functions. (c_omp_check_loop_iv): Handle generated loops. For now sorry on mixing non-rectangular loop with generated loops. (c_omp_check_loop_binding_exprs): For now sorry on mixing imperfect loops with generated loops. (c_omp_directives): Uncomment tile and unroll entries. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL, change PRAGMA_OMP__LAST_ to the latter. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_FULL and PRAGMA_OMP_CLAUSE_PARTIAL. * c-pragma.cc (omp_pragmas_simd): Add tile and unroll omp pragmas. gcc/c/ * c-parser.cc (c_parser_skip_std_attribute_spec_seq): New function. (check_omp_intervening_code): Reject imperfectly nested tile. (c_parser_compound_statement_nostart): If want_nested_loop, use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. (c_parser_omp_clause_name): Handle full and partial clause names. (c_parser_omp_clause_allocate): Remove spurious semicolon. (c_parser_omp_clause_full, c_parser_omp_clause_partial): New functions. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_FULL and PRAGMA_OMP_CLAUSE_PARTIAL. (c_parser_omp_next_tokens_can_be_canon_loop): New function. (c_parser_omp_loop_nest): Parse C23 attributes. Handle tile/unroll constructs. Use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Only add_stmt (body) if it is non-NULL. (c_parser_omp_for_loop): Rename tiling variable to oacc_tiling. For OMP_CLAUSE_SIZES set collapse to list length of OMP_CLAUSE_SIZES_LIST. Use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Remove spurious semicolon. Don't call c_omp_check_loop_binding_exprs if stmt is NULL. Skip generated loops. (c_parser_omp_tile_sizes, c_parser_omp_tile): New functions. (OMP_UNROLL_CLAUSE_MASK): Define. (c_parser_omp_unroll): New function. (c_parser_omp_construct): Handle PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL. * c-typeck.cc (c_finish_omp_clauses): Adjust wording of some of the conflicting clause diagnostic messages to include word clause. Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES} and diagnose full vs. partial conflict. gcc/cp/ * cp-tree.h (dependent_omp_for_p): Add another tree argument. * parser.cc (check_omp_intervening_code): Reject imperfectly nested tile. (cp_parser_statement_seq_opt): If want_nested_loop, use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. (cp_parser_omp_clause_name): Handle full and partial clause names. (cp_parser_omp_clause_full, cp_parser_omp_clause_partial): New functions. (cp_parser_omp_all_clauses): Formatting fix. Handle PRAGMA_OMP_CLAUSE_PARTIAL and PRAGMA_OMP_CLAUSE_FULL. (cp_parser_next_tokens_can_be_canon_loop): New function. (cp_parser_omp_loop_nest): Parse C++11 attributes. Handle tile/unroll constructs. Use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Only add_stmt cp_parser_omp_loop_nest result if it is non-NULL. (cp_parser_omp_for_loop): Rename tiling variable to oacc_tiling. For OMP_CLAUSE_SIZES set collapse to list length of OMP_CLAUSE_SIZES_LIST. Use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Remove spurious semicolon. Don't call c_omp_check_loop_binding_exprs if stmt is NULL. Skip and/or handle generated loops. Remove spurious ()s around & operands. (cp_parser_omp_tile_sizes, cp_parser_omp_tile): New functions. (OMP_UNROLL_CLAUSE_MASK): Define. (cp_parser_omp_unroll): New function. (cp_parser_omp_construct): Handle PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL. (cp_parser_pragma): Likewise. * semantics.cc (finish_omp_clauses): Don't call fold_build_cleanup_point_expr for cases which obviously won't need it, like checked INTEGER_CSTs. Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES} and diagnose full vs. partial conflict. Adjust wording of some of the conflicting clause diagnostic messages to include word clause. (finish_omp_for): Use decl equal to global_namespace as a marker for generated loop. Pass also body to dependent_omp_for_p. Skip generated loops. (finish_omp_for_block): Skip generated loops. * pt.cc (tsubst_omp_clauses): Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES}. (tsubst_stmt): Handle OMP_TILE and OMP_UNROLL. Handle or skip generated loops. (dependent_omp_for_p): Add body argument. If declv vector element is NULL, find generated loop. * cp-gimplify.cc (cp_gimplify_expr): Handle OMP_TILE and OMP_UNROLL. (cp_fold_r): Likewise. (cp_genericize_r): Likewise. Skip generated loops. gcc/fortran/ * gfortran.h (enum gfc_statement): Add ST_OMP_UNROLL, ST_OMP_END_UNROLL, ST_OMP_TILE and ST_OMP_END_TILE. (struct gfc_omp_clauses): Add sizes_list, partial, full and erroneous members. (enum gfc_exec_op): Add EXEC_OMP_UNROLL and EXEC_OMP_TILE. (gfc_expr_list_len): Declare. * match.h (gfc_match_omp_tile, gfc_match_omp_unroll): Declare. * openmp.cc (gfc_get_location): Declare. (gfc_free_omp_clauses): Free sizes_list. (match_oacc_expr_list): Rename to ... (match_omp_oacc_expr_list): ... this. Add is_omp argument and change diagnostic wording if it is true. (enum omp_mask2): Add OMP_CLAUSE_{FULL,PARTIAL,SIZES}. (gfc_match_omp_clauses): Parse full, partial and sizes clauses. (gfc_match_oacc_wait): Use match_omp_oacc_expr_list instead of match_oacc_expr_list. (OMP_UNROLL_CLAUSES, OMP_TILE_CLAUSES): Define. (gfc_match_omp_tile, gfc_match_omp_unroll): New functions. (resolve_omp_clauses): Diagnose full vs. partial clause conflict. Resolve sizes clause arguments. (find_nested_loop_in_chain): Use switch instead of series of ifs. Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_omp_do_blocks): Set omp_current_do_collapse to list length of sizes_list if present. (gfc_resolve_do_iterator): Return for EXEC_OMP_TILE or EXEC_OMP_UNROLL. (restructure_intervening_code): Remove spurious ()s around & operands. (is_outer_iteration_variable): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (check_nested_loop_in_chain): Likewise. (expr_is_invariant): Likewise. (resolve_omp_do): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. Diagnose tile without sizes clause. Use sizes_list length for count if non-NULL. Set code->ext.omp_clauses->erroneous on loops where we've reported diagnostics. Sorry for mixing non-rectangular loops with generated loops. (omp_code_to_statement): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_omp_directive): Likewise. * parse.cc (decode_omp_directive): Parse end tile, end unroll, tile and unroll. Move nothing entry alphabetically. (case_exec_markers): Add ST_OMP_TILE and ST_OMP_UNROLL. (gfc_ascii_statement): Handle ST_OMP_END_TILE, ST_OMP_END_UNROLL, ST_OMP_TILE and ST_OMP_UNROLL. (parse_omp_do): Add nested argument. Handle ST_OMP_TILE and ST_OMP_UNROLL. (parse_omp_structured_block): Adjust parse_omp_do caller. (parse_executable): Likewise. Handle ST_OMP_TILE and ST_OMP_UNROLL. * resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_code): Likewise. * st.cc (gfc_free_statement): Likewise. * trans.cc (trans_code): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses): Handle full, partial and sizes clauses. Use tree_cons + nreverse instead of temporary vector and build_tree_list_vec for tile_list handling. (gfc_expr_list_len): New function. (gfc_trans_omp_do): Rename tile to oacc_tile. Handle sizes clause. Don't assert code->op is EXEC_DO. Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_trans_omp_directive): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. * dump-parse-tree.cc (show_omp_clauses): Dump full, partial and sizes clauses. (show_omp_node): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (show_code_node): Likewise. gcc/testsuite/ * c-c++-common/gomp/attrs-tile-1.c: New test. * c-c++-common/gomp/attrs-tile-2.c: New test. * c-c++-common/gomp/attrs-tile-3.c: New test. * c-c++-common/gomp/attrs-tile-4.c: New test. * c-c++-common/gomp/attrs-tile-5.c: New test. * c-c++-common/gomp/attrs-tile-6.c: New test. * c-c++-common/gomp/attrs-unroll-1.c: New test. * c-c++-common/gomp/attrs-unroll-2.c: New test. * c-c++-common/gomp/attrs-unroll-3.c: New test. * c-c++-common/gomp/attrs-unroll-inner-1.c: New test. * c-c++-common/gomp/attrs-unroll-inner-2.c: New test. * c-c++-common/gomp/attrs-unroll-inner-3.c: New test. * c-c++-common/gomp/attrs-unroll-inner-4.c: New test. * c-c++-common/gomp/attrs-unroll-inner-5.c: New test. * c-c++-common/gomp/imperfect-attributes.c: Adjust expected diagnostics. * c-c++-common/gomp/imperfect-loop-nest.c: New test. * c-c++-common/gomp/ordered-5.c: New test. * c-c++-common/gomp/scan-7.c: New test. * c-c++-common/gomp/tile-1.c: New test. * c-c++-common/gomp/tile-2.c: New test. * c-c++-common/gomp/tile-3.c: New test. * c-c++-common/gomp/tile-4.c: New test. * c-c++-common/gomp/tile-5.c: New test. * c-c++-common/gomp/tile-6.c: New test. * c-c++-common/gomp/tile-7.c: New test. * c-c++-common/gomp/tile-8.c: New test. * c-c++-common/gomp/tile-9.c: New test. * c-c++-common/gomp/tile-10.c: New test. * c-c++-common/gomp/tile-11.c: New test. * c-c++-common/gomp/tile-12.c: New test. * c-c++-common/gomp/tile-13.c: New test. * c-c++-common/gomp/tile-14.c: New test. * c-c++-common/gomp/tile-15.c: New test. * c-c++-common/gomp/unroll-1.c: New test. * c-c++-common/gomp/unroll-2.c: New test. * c-c++-common/gomp/unroll-3.c: New test. * c-c++-common/gomp/unroll-4.c: New test. * c-c++-common/gomp/unroll-5.c: New test. * c-c++-common/gomp/unroll-6.c: New test. * c-c++-common/gomp/unroll-7.c: New test. * c-c++-common/gomp/unroll-8.c: New test. * c-c++-common/gomp/unroll-9.c: New test. * c-c++-common/gomp/unroll-inner-1.c: New test. * c-c++-common/gomp/unroll-inner-2.c: New test. * c-c++-common/gomp/unroll-inner-3.c: New test. * c-c++-common/gomp/unroll-non-rect-1.c: New test. * c-c++-common/gomp/unroll-non-rect-2.c: New test. * c-c++-common/gomp/unroll-non-rect-3.c: New test. * c-c++-common/gomp/unroll-simd-1.c: New test. * gcc.dg/gomp/attrs-4.c: Adjust expected diagnostics. * gcc.dg/gomp/for-1.c: Likewise. * gcc.dg/gomp/for-11.c: Likewise. * g++.dg/gomp/attrs-4.C: Likewise. * g++.dg/gomp/for-1.C: Likewise. * g++.dg/gomp/pr94512.C: Likewise. * g++.dg/gomp/tile-1.C: New test. * g++.dg/gomp/tile-2.C: New test. * g++.dg/gomp/unroll-1.C: New test. * g++.dg/gomp/unroll-2.C: New test. * g++.dg/gomp/unroll-3.C: New test. * gfortran.dg/gomp/inner-loops-1.f90: New test. * gfortran.dg/gomp/inner-loops-2.f90: New test. * gfortran.dg/gomp/pure-1.f90: Add tests for !$omp unroll and !$omp tile. * gfortran.dg/gomp/pure-2.f90: Remove those tests from here. * gfortran.dg/gomp/scan-9.f90: New test. * gfortran.dg/gomp/tile-1.f90: New test. * gfortran.dg/gomp/tile-2.f90: New test. * gfortran.dg/gomp/tile-3.f90: New test. * gfortran.dg/gomp/tile-4.f90: New test. * gfortran.dg/gomp/tile-5.f90: New test. * gfortran.dg/gomp/tile-6.f90: New test. * gfortran.dg/gomp/tile-7.f90: New test. * gfortran.dg/gomp/tile-8.f90: New test. * gfortran.dg/gomp/tile-9.f90: New test. * gfortran.dg/gomp/tile-10.f90: New test. * gfortran.dg/gomp/tile-imperfect-nest-1.f90: New test. * gfortran.dg/gomp/tile-imperfect-nest-2.f90: New test. * gfortran.dg/gomp/tile-inner-loops-1.f90: New test. * gfortran.dg/gomp/tile-inner-loops-2.f90: New test. * gfortran.dg/gomp/tile-inner-loops-3.f90: New test. * gfortran.dg/gomp/tile-inner-loops-4.f90: New test. * gfortran.dg/gomp/tile-inner-loops-5.f90: New test. * gfortran.dg/gomp/tile-inner-loops-6.f90: New test. * gfortran.dg/gomp/tile-inner-loops-7.f90: New test. * gfortran.dg/gomp/tile-inner-loops-8.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-1.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-2.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-3.f90: New test. * gfortran.dg/gomp/tile-unroll-1.f90: New test. * gfortran.dg/gomp/tile-unroll-2.f90: New test. * gfortran.dg/gomp/unroll-1.f90: New test. * gfortran.dg/gomp/unroll-2.f90: New test. * gfortran.dg/gomp/unroll-3.f90: New test. * gfortran.dg/gomp/unroll-4.f90: New test. * gfortran.dg/gomp/unroll-5.f90: New test. * gfortran.dg/gomp/unroll-6.f90: New test. * gfortran.dg/gomp/unroll-7.f90: New test. * gfortran.dg/gomp/unroll-8.f90: New test. * gfortran.dg/gomp/unroll-9.f90: New test. * gfortran.dg/gomp/unroll-10.f90: New test. * gfortran.dg/gomp/unroll-11.f90: New test. * gfortran.dg/gomp/unroll-12.f90: New test. * gfortran.dg/gomp/unroll-13.f90: New test. * gfortran.dg/gomp/unroll-inner-loop-1.f90: New test. * gfortran.dg/gomp/unroll-inner-loop-2.f90: New test. * gfortran.dg/gomp/unroll-no-clause-1.f90: New test. * gfortran.dg/gomp/unroll-non-rect-1.f90: New test. * gfortran.dg/gomp/unroll-non-rect-2.f90: New test. * gfortran.dg/gomp/unroll-simd-1.f90: New test. * gfortran.dg/gomp/unroll-simd-2.f90: New test. * gfortran.dg/gomp/unroll-simd-3.f90: New test. * gfortran.dg/gomp/unroll-tile-1.f90: New test. * gfortran.dg/gomp/unroll-tile-2.f90: New test. * gfortran.dg/gomp/unroll-tile-inner-1.f90: New test. libgomp/ * testsuite/libgomp.c-c++-common/imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/matrix-1.h: New test. * testsuite/libgomp.c-c++-common/matrix-constant-iter.h: New test. * testsuite/libgomp.c-c++-common/matrix-helper.h: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-unroll-full-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-transform-variants-1.h: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/unroll-1.c: New test. * testsuite/libgomp.c-c++-common/unroll-non-rect-1.c: New test. * testsuite/libgomp.c++/matrix-no-directive-unroll-full-1.C: New test. * testsuite/libgomp.c++/tile-2.C: New test. * testsuite/libgomp.c++/tile-3.C: New test. * testsuite/libgomp.c++/unroll-1.C: New test. * testsuite/libgomp.c++/unroll-2.C: New test. * testsuite/libgomp.c++/unroll-full-tile.C: New test. * testsuite/libgomp.fortran/imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/inner-1.f90: New test. * testsuite/libgomp.fortran/nested-fn.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/tile-1.f90: New test. * testsuite/libgomp.fortran/tile-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-1.f90: New test. * testsuite/libgomp.fortran/tile-unroll-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-3.f90: New test. * testsuite/libgomp.fortran/tile-unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-1.f90: New test. * testsuite/libgomp.fortran/unroll-2.f90: New test. * testsuite/libgomp.fortran/unroll-3.f90: New test. * testsuite/libgomp.fortran/unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-5.f90: New test. * testsuite/libgomp.fortran/unroll-6.f90: New test. * testsuite/libgomp.fortran/unroll-7a.f90: New test. * testsuite/libgomp.fortran/unroll-7b.f90: New test. * testsuite/libgomp.fortran/unroll-7c.f90: New test. * testsuite/libgomp.fortran/unroll-7.f90: New test. * testsuite/libgomp.fortran/unroll-8.f90: New test. * testsuite/libgomp.fortran/unroll-simd-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-2.f90: New test. |
||
|
Rainer Orth
|
7e322d576e |
testsuite: Adjust several dg-additional-files-options calls [PR115294]
A recent patch
commit
|
||
|
Jakub Jelinek
|
5eb25d1561 |
libgomp: Add gfx90c, 1036 and 1103 declare variant tests
Recently -march=gfx{90c,1036,1103} support has been added, but corresponding
changes weren't done in the testsuite.
The following patch adds that.
Tested on x86_64-linux (with fiji and gfx1103 devices; had to use
OMP_DEFAULT_DEVICE=1 there, fiji doesn't really work due to LLVM dropping
support, but we still list those as offloading devices).
2024-05-02 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/declare-variant-4.h (gfx90c, gfx1036, gfx1103):
New functions.
(f): Add #pragma omp declare variant directives for those.
* testsuite/libgomp.c/declare-variant-4-gfx90c.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1036.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1103.c: New test.
|
||
|
Chung-Lin Tang
|
a7578a077e |
OpenACC 2.7: Adjust acc_map_data/acc_unmap_data interaction with reference counters
This patch adjusts the implementation of acc_map_data/acc_unmap_data API library routines to more fit the description in the OpenACC 2.7 specification. Instead of using REFCOUNT_INFINITY, we now define a REFCOUNT_ACC_MAP_DATA special value to mark acc_map_data-created mappings. Adjustment around mapping related code to respect OpenACC semantics are also added. libgomp/ChangeLog: * libgomp.h (REFCOUNT_ACC_MAP_DATA): Define as (REFCOUNT_SPECIAL | 2). * oacc-mem.c (acc_map_data): Adjust to use REFCOUNT_ACC_MAP_DATA, initialize dynamic_refcount as 1. (acc_unmap_data): Adjust to use REFCOUNT_ACC_MAP_DATA, (goacc_map_var_existing): Add REFCOUNT_ACC_MAP_DATA case. (goacc_exit_datum_1): Add REFCOUNT_ACC_MAP_DATA case, respect REFCOUNT_ACC_MAP_DATA when decrementing/finalizing. Force lowest dynamic_refcount to be 1 for REFCOUNT_ACC_MAP_DATA. (goacc_enter_data_internal): Add REFCOUNT_ACC_MAP_DATA case. * target.c (gomp_increment_refcount): Return early for REFCOUNT_ACC_MAP_DATA case. (gomp_decrement_refcount): Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-96.c: New testcase. * testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c: Adjust testcase error output scan test. |
||
|
Thomas Schwinge
|
679f81a32f |
nvptx: In mkoffload.cc, call diagnostic_color_init + gcc_init_libintl: Restore 'libgomp.c/reverse-offload-sm30.c' testing
With commit
|
||
|
Jakub Jelinek
|
592536eb3c |
c++: Fix ICE with weird copy assignment operator [PR114572]
While ctors/dtors don't return anything (undeclared void or this pointer on arm) and copy assignment operators normally return a reference to *this, it isn't invalid to return uselessly some class object which might need destructing, but the OpenMP clause handling code wasn't expecting that. The following patch fixes that. 2024-04-05 Jakub Jelinek <jakub@redhat.com> PR c++/114572 * cp-gimplify.cc (cxx_omp_clause_apply_fn): Call build_cplus_new on build_call_a result if it has class type. * testsuite/libgomp.c++/pr114572.C: New test. |
||
|
Kwok Cheung Yeung
|
637e76b90e |
openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls
A splay-tree was previously used to lookup equivalent target addresses for a given host address on offload targets. However, as splay-trees can modify their structure on lookup, they are not suitable for concurrent access from separate teams/threads without some form of locking. This patch changes the lookup data structure to a hashtab instead, which does not have these issues. The call to build_indirect_map to initialize the data structure is now called from just the first thread of the first team to avoid redundant calls to this function. 2024-03-22 Kwok Cheung Yeung <kcyeung@baylibre.com> libgomp/ * config/accel/target-indirect.c: Include string.h and hashtab.h. Remove include of splay-tree.h. Update comments. (splay_tree_prefix, splay_tree_c): Delete. (struct indirect_map_t): New. (hash_entry_type, htab_alloc, htab_free, htab_hash, htab_eq): New. (GOMP_INDIRECT_ADD_MAP): Remove volatile qualifier. (USE_SPLAY_TREE_LOOKUP): Rename to... (USE_HASHTAB_LOOKUP): ..this. (indirect_map, indirect_array): Delete. (indirect_htab): New. (build_indirect_map): Remove locking. Build indirect map using hashtab. (GOMP_target_map_indirect_ptr): Use indirect_htab to lookup target address. (GOMP_target_map_indirect_ptr): Remove volatile qualifier. * config/gcn/team.c (gomp_gcn_enter_kernel): Call build_indirect_map from first thread of first team only. * config/nvptx/team.c (gomp_nvptx_main): Likewise. * testsuite/libgomp.c-c++-common/declare-target-indirect-2.c (main): Add missing break statements. * testsuite/libgomp.fortran/declare-target-indirect-2.f90: Remove xfail. |
||
|
Thomas Schwinge
|
25242ed8eb |
Fix 'char' initialization, copy, check in 'libgomp.oacc-fortran/acc-memcpy.f90'
Our dear friend '-Wuninitialized' reported:
[...]/libgomp.oacc-fortran/acc-memcpy.f90:18:27:
18 | char(j) = int (j, int8)
| ^
Warning: ‘j’ may be used uninitialized [-Wmaybe-uninitialized]
[...]/libgomp.oacc-fortran/acc-memcpy.f90:14:20:
14 | integer(int8) :: j
| ^
note: ‘j’ was declared here
..., but actually there were other issues.
libgomp/
* testsuite/libgomp.oacc-fortran/acc-memcpy.f90: Fix 'char'
initialization, copy, check.
|
||
|
Tobias Burnus
|
c5037fcee2 |
OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]
Dummy procedures look similar to variables but aren't - neither in Fortran nor in OpenMP. As the middle end sees PARM_DECLs, mark them as predetermined firstprivate for mapping (as already done in gfc_omp_predetermined_sharing). This does not address the isses related to procedure pointers, which are still discussed on spec level [see PR]. PR fortran/114283 gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_omp_predetermined_mapping): Map dummy procedures as firstprivate. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-indirect-4.f90: New test. |
||
|
John David Anglin
|
49c3f24552 |
Revert "Set num_threads to 50 on 32-bit hppa in two libgomp loop tests"
This reverts commit
|
||
|
Jakub Jelinek
|
4f82d5a95a |
OpenMP/C++: Fix (first)private clause with member variables [PR110347]
OpenMP permits '(first)private' for C++ member variables, which GCC handles by tagging those by DECL_OMP_PRIVATIZED_MEMBER, adding a temporary VAR_DECL and DECL_VALUE_EXPR pointing to the 'this->member_var' in the C++ front end. The idea is that in omp-low.cc, the DECL_VALUE_EXPR is used before the region (for 'firstprivate'; ignored for 'private') while in the region, the DECL itself is used. In gimplify, the value expansion is suppressed and deferred if the lang_hooks.decls.omp_disregard_value_expr (decl, shared) returns true - which is never the case if 'shared' is true. In OpenMP 4.5, only 'map' and 'use_device_ptr' was permitted for the 'target' directive. And when OpenMP 5.0's 'private'/'firstprivate' clauses was added, the the update that now 'shared' argument could be false was missed. The respective check has now been added. 2024-03-01 Jakub Jelinek <jakub@redhat.com> Tobias Burnus <tburnus@baylibre.com> PR c++/110347 gcc/ChangeLog: * gimplify.cc (omp_notice_variable): Fix 'shared' arg to lang_hooks.decls.omp_disregard_value_expr for (first)private in target regions. libgomp/ChangeLog: * testsuite/libgomp.c++/target-lambda-3.C: Moved from gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling. * testsuite/libgomp.c++/target-lambda-1.C: Modify to also also work without offloading. * testsuite/libgomp.c++/firstprivate-1.C: New test. * testsuite/libgomp.c++/firstprivate-2.C: New test. * testsuite/libgomp.c++/private-1.C: New test. * testsuite/libgomp.c++/private-2.C: New test. * testsuite/libgomp.c++/target-lambda-4.C: New test. * testsuite/libgomp.c++/use_device_ptr-1.C: New test. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-lambda-1.C: Moved to become a run-time test under testsuite/libgomp.c++. Co-authored-by: Tobias Burnus <tburnus@baylibre.com> |
||
|
Tobias Burnus
|
8b3f1edf9b |
OpenACC: Add Fortran routines acc_{alloc,free,hostptr,deviceptr,memcpy_{to,from}_device*}
These routines map simply to the C counterpart and are meanwhile defined in OpenACC 3.3. (There are additional routine changes, including the Fortran addition of acc_attach/acc_detach, that require more work than a simple addition of an interface and are therefore excluded.) libgomp/ChangeLog: * libgomp.texi (OpenACC Runtime Library Routines): Document new 3.3 routines that simply map to their C counterpart. * openacc.f90 (openacc): Add them. * openacc_lib.h: Likewise. * testsuite/libgomp.oacc-fortran/acc_host_device_ptr.f90: New test. * testsuite/libgomp.oacc-fortran/acc-memcpy.f90: New test. * testsuite/libgomp.oacc-fortran/acc-memcpy-2.f90: New test. * testsuite/libgomp.oacc-c-c++-common/lib-59.c: Crossref to f90 test. * testsuite/libgomp.oacc-c-c++-common/lib-60.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-95.c: Likewise. |
||
|
Kwok Cheung Yeung
|
451bb58660 |
openmp, fortran: Add Fortran support for indirect clause on the declare target directive
2024-02-15 Kwok Cheung Yeung <kcyeung@baylibre.com> gcc/fortran/ * dump-parse-tree.cc (show_attr): Handle omp_declare_target_indirect attribute. * f95-lang.cc (gfc_gnu_attributes): Add entry for 'omp declare target indirect'. * gfortran.h (symbol_attribute): Add omp_declare_target_indirect field. (struct gfc_omp_clauses): Add indirect field. * openmp.cc (omp_mask2): Add OMP_CLAUSE_INDIRECT. (gfc_match_omp_clauses): Match indirect clause. (OMP_DECLARE_TARGET_CLAUSES): Add OMP_CLAUSE_INDIRECT. (gfc_match_omp_declare_target): Check omp_device_type and apply omp_declare_target_indirect attribute to symbol if indirect clause active. Show warning if there are only device_type and/or indirect clauses on the directive. * trans-decl.cc (add_attributes_to_decl): Add 'omp declare target indirect' attribute if symbol has indirect attribute set. gcc/testsuite/ * gfortran.dg/gomp/declare-target-4.f90 (f1): Update expected warning. * gfortran.dg/gomp/declare-target-indirect-1.f90: New. * gfortran.dg/gomp/declare-target-indirect-2.f90: New. libgomp/ * testsuite/libgomp.fortran/declare-target-indirect-1.f90: New. * testsuite/libgomp.fortran/declare-target-indirect-2.f90: New. * testsuite/libgomp.fortran/declare-target-indirect-3.f90: New. |
||
|
Rainer Orth
|
1e94648ab7 |
libgomp: testsuite: Don't XPASS libgomp.c/alloc-pinned-1.c etc. on non-Linux targets [PR113448]
Two libgomp tests XPASS on Solaris (any non-Linux target actually) since
their introduction:
XPASS: libgomp.c/alloc-pinned-1.c execution test
XPASS: libgomp.c/alloc-pinned-2.c execution test
The problem is that the test just prints
OS unsupported
and exits successfully, while the test is XFAILed:
/* { dg-xfail-run-if "Pinning not implemented on this host" { ! *-*-linux-gnu } } */
Fixed by aborting immediately after the message above in the non-Linux
case.
Tested on i386-pc-solaris2.11 and i686-pc-linux-gnu.
2024-02-02 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE>
libgomp:
PR testsuite/113448
* testsuite/libgomp.c/alloc-pinned-1.c [!__linux__] (CHECK_SIZE):
Call abort.
* testsuite/libgomp.c/alloc-pinned-2.c [!__linux__] (CHECK_SIZE):
Likewise.
|
||
|
John David Anglin
|
b14209715e |
Set num_threads to 50 on 32-bit hppa in two libgomp loop tests
We support a maximum of 50 threads on 32-bit hppa. 2024-02-01 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.c++/loop-3.C: Set num_threads to 50 on 32-bit hppa. * testsuite/libgomp.c/omp-loop03.c: Likewise. |
||
|
Tobias Burnus
|
cb366731e7 |
libgomp.c/declare-variant-4.h: Fix used variant function for gfx1030/gfx1100
libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4.h: Use gfx1100/gfx1030 function not gfx90a for gfx1100/gfx1030 context selector. Signed-off-by: Tobias Burnus <tburnus@baylibre.com> |
||
|
Tobias Burnus
|
56d0aba11a |
amdgcn: config.gcc - enable gfx1030 and gfx1100 multilib; add them to the docs
gcc/ChangeLog: * config.gcc (amdgcn-*-*): Add gfx1030 and gfx1100 to TM_MULTILIB_CONFIG. * doc/install.texi (Configuration amdgcn-*-*): Mention gfx1030/gfx1100. * doc/invoke.texi (AMD GCN Options): Add gfx1030 and gfx1100 to -march/-mtune. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4.h: Add variant functions for gfx1030 and gfx1100. * testsuite/libgomp.c/declare-variant-4-gfx1030.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx1100.c: New test. Signed-off-by: Tobias Burnus <tburnus@baylibre.com> |
||
|
Tobias Burnus
|
a1b2953924 |
xfail libgomp.c/declare-variant-4-{fiji,gfx803}.c
Since r14-4734-g56ed1055b2f40ac162ae8d382280ac07a33f789f, GCC no longer builds the Fiji (alias gfx803) libraries by default as support for it was removed in ROCm 4.0 and will be removed in LLVM 18. Thus, unless gfx803 is explicitly enabled, the following testcases will fail to link as libgomp is not available for Fiji. Hence, this commit xfails those testcases. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4-fiji.c: Xfail as fiji support is no longer enabled by default. * testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise. Signed-off-by: Tobias Burnus <tburnus@baylibre.com> |
||
|
John David Anglin
|
c9c507ff8b |
Increase timeout by 2 in libgomp.fortran/alloc-comp-3.f90 on hppa*-*-*
2024-01-20 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.fortran/alloc-comp-3.f90: Increase timeout by 2 on hppa*-*-*. |
||
|
John David Anglin
|
584e43e434 |
Don't run libgomp.c/simd-math-1.c on hppa*-*-hpux*
hppa*-*-hpux* lacks necessary math functions. 2024-01-20 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.c/simd-math-1.c: Don't run on hppa*-*-hpux*. |
||
|
Jakub Jelinek
|
c8f1045679 |
openmp: Add OpenMP _BitInt support [PR113409]
The following patch adds support for _BitInt iterators of OpenMP canonical loops (with the preexisting limitation that when not using compile time static scheduling the iterators in the library are at most unsigned long long or signed long, so one can't in the runtime/dynamic/guided etc. cases iterate more than what those types can represent, like is the case of e.g. __int128 iterators too) and the testcase also covers linear/reduction clauses for them. 2024-01-17 Jakub Jelinek <jakub@redhat.com> PR middle-end/113409 * omp-general.cc (omp_adjust_for_condition): Handle BITINT_TYPE like INTEGER_TYPE. (omp_extract_for_data): Use build_bitint_type rather than build_nonstandard_integer_type if either iter_type or loop->v type is BITINT_TYPE. * omp-expand.cc (expand_omp_for_generic, expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): Handle BITINT_TYPE like INTEGER_TYPE. * testsuite/libgomp.c/bitint-1.c: New test. |
||
|
Julian Brown
|
b5476e4c88 |
OpenMP: lvalue parsing for map/to/from clauses (C)
This patch adds support for parsing general lvalues ("locator list item
types") for OpenMP "map", "to" and "from" clauses to the C front-end,
similar to the previously-posted patch for C++. Such syntax is permitted
for OpenMP 5.0 and above. It was previously posted for mainline here:
https://gcc.gnu.org/pipermail/gcc-patches/2022-December/609038.html
and for the og13 branch here:
https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623355.html
2024-01-11 Julian Brown <julian@codesourcery.com>
gcc/c-family/
* c-pretty-print.cc (c_pretty_printer::postfix_expression,
c_pretty_printer::expression): Add OMP_ARRAY_SECTION support.
gcc/c/
* c-parser.cc (c_parser_braced_init, c_parser_conditional_expression):
Don't allow OpenMP array section.
(c_parser_postfix_expression): Don't allow array section in statement
expression.
(c_parser_postfix_expression_after_primary): Add support for OpenMP
array section parsing.
(c_parser_expr_list): Don't allow OpenMP array section here.
(c_parser_omp_variable_list): Change ALLOW_DEREF parameter to
MAP_LVALUE. Support parsing of general lvalues in "map", "to" and
"from" clauses.
(c_parser_omp_var_list_parens): Change ALLOW_DEREF parameter to
MAP_LVALUE. Update call to c_parser_omp_variable_list.
(c_parser_oacc_data_clause): Update calls to
c_parser_omp_var_list_parens.
(c_parser_omp_clause_reduction): Use OMP_ARRAY_SECTION tree node
instead of TREE_LIST for array sections.
(c_parser_omp_target): Allow GOMP_MAP_ATTACH.
* c-tree.h (c_omp_array_section_p): Add extern declaration.
(build_omp_array_section): Add prototype.
* c-typeck.cc (c_omp_array_section_p): Add flag.
(mark_exp_read): Support OMP_ARRAY_SECTION.
(build_omp_array_section): Add function.
(build_external_ref): Tweak error path for OpenMP array sections.
(handle_omp_array_sections_1): Use OMP_ARRAY_SECTION tree code instead
of TREE_LIST. Handle more kinds of expressions.
(c_oacc_check_attachments): Use OMP_ARRAY_SECTION instead of TREE_LIST
for array sections.
(c_finish_omp_clauses): Use OMP_ARRAY_SECTION instead of TREE_LIST.
Check for supported expression types.
gcc/testsuite/
* gcc.dg/gomp/bad-array-section-c-1.c: New test.
* gcc.dg/gomp/bad-array-section-c-2.c: New test.
* gcc.dg/gomp/bad-array-section-c-3.c: New test.
* gcc.dg/gomp/bad-array-section-c-4.c: New test.
* gcc.dg/gomp/bad-array-section-c-5.c: New test.
* gcc.dg/gomp/bad-array-section-c-6.c: New test.
* gcc.dg/gomp/bad-array-section-c-7.c: New test.
* gcc.dg/gomp/bad-array-section-c-8.c: New test.
libgomp/
* libgomp.texi: C/C++ lvalues are supported now for map/to/from.
* testsuite/libgomp.c-c++-common/ind-base-4.c: New test.
* testsuite/libgomp.c-c++-common/unary-ptr-1.c: New test.
|
||
|
Julian Brown
|
1413af02d6 |
OpenMP: lvalue parsing for map/to/from clauses (C++)
This patch supports "lvalue" parsing (or "locator list item type" parsing) for several OpenMP clause types for C++, as required for OpenMP 5.0 and above. This version has been rebased -- some things have changed around template handling recently, e.g. removal of build_non_dependent_expr and tsubst_copy. A new potential corner-case issue has shown up regarding implicit mapping of references to pointer to pointers -- an interaction with the post-review fixes/rework for the patch here: https://gcc.gnu.org/pipermail/gcc-patches/2023-November/638602.html Which fixed the (new) tests baseptrs-[6789].C. I've noted that for now in the patch, and adjusted the baseptrs-[46].C tests slightly to accommodate. 2024-01-08 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-common.h (c_omp_address_inspector): Remove static from get_origin and maybe_unconvert_ref methods. * c-omp.cc (c_omp_split_clauses): Support OMP_ARRAY_SECTION. (c_omp_address_inspector::map_supported_p): Handle OMP_ARRAY_SECTION. (c_omp_address_inspector::get_origin): Avoid dereferencing possibly NULL type when processing template decls. (c_omp_address_inspector::maybe_unconvert_ref): Likewise. gcc/cp/ * constexpr.cc (potential_consant_expression_1): Handle OMP_ARRAY_SECTION. * cp-tree.h (grok_omp_array_section, build_omp_array_section): Add prototypes. * decl2.cc (grok_omp_array_section): New function. * error.cc (dump_expr): Handle OMP_ARRAY_SECTION. * parser.cc (cp_parser_new): Initialize parser->omp_array_section_p. (cp_parser_statement_expr): Disallow array sections. (cp_parser_postfix_open_square_expression): Support OMP_ARRAY_SECTION parsing. (cp_parser_parenthesized_expression_list, cp_parser_lambda_expression, cp_parser_braced_list): Disallow array sections. (cp_parser_omp_var_list_no_open): Remove ALLOW_DEREF parameter, add MAP_LVALUE in its place. Support generalised lvalue parsing for OpenMP map, to and from clauses. Use OMP_ARRAY_SECTION code instead of TREE_LIST to represent OpenMP array sections. (cp_parser_omp_var_list): Remove ALLOW_DEREF parameter, add MAP_LVALUE. Pass to cp_parser_omp_var_list_no_open. (cp_parser_oacc_data_clause): Update call to cp_parser_omp_var_list. (cp_parser_omp_clause_map): Add sk_omp scope around cp_parser_omp_var_list_no_open call. * parser.h (cp_parser): Add omp_array_section_p field. * pt.cc (tsubst, tsubst_copy, tsubst_omp_clause_decl, tsubst_copy_and_build): Add OMP_ARRAY_SECTION support. * semantics.cc (handle_omp_array_sections_1, handle_omp_array_sections, cp_oacc_check_attachments, finish_omp_clauses): Use OMP_ARRAY_SECTION instead of TREE_LIST where appropriate. Handle more types of map expression. * typeck.cc (build_omp_array_section): New function. gcc/ * gimplify.cc (gimplify_expr): Ensure OMP_ARRAY_SECTION has been processed out before gimplification. * tree-pretty-print.cc (dump_generic_node): Support OMP_ARRAY_SECTION. * tree.def (OMP_ARRAY_SECTION): New tree code. gcc/testsuite/ * c-c++-common/gomp/map-6.c: Update expected output. * c-c++-common/gomp/target-enter-data-1.c: Update scan test. * g++.dg/gomp/array-section-1.C: New test. * g++.dg/gomp/array-section-2.C: New test. * g++.dg/gomp/bad-array-section-1.C: New test. * g++.dg/gomp/bad-array-section-2.C: New test. * g++.dg/gomp/bad-array-section-3.C: New test. * g++.dg/gomp/bad-array-section-4.C: New test. * g++.dg/gomp/bad-array-section-5.C: New test. * g++.dg/gomp/bad-array-section-6.C: New test. * g++.dg/gomp/bad-array-section-7.C: New test. * g++.dg/gomp/bad-array-section-8.C: New test. * g++.dg/gomp/bad-array-section-9.C: New test. * g++.dg/gomp/bad-array-section-10.C: New test. * g++.dg/gomp/bad-array-section-11.C: New test. * g++.dg/gomp/has_device_addr-non-lvalue-1.C: New test. * g++.dg/gomp/pr67522.C: Update expected output. * g++.dg/gomp/ind-base-3.C: New test. * g++.dg/gomp/map-assignment-1.C: New test. * g++.dg/gomp/map-inc-1.C: New test. * g++.dg/gomp/map-lvalue-ref-1.C: New test. * g++.dg/gomp/map-ptrmem-1.C: New test. * g++.dg/gomp/map-ptrmem-2.C: New test. * g++.dg/gomp/map-static-cast-lvalue-1.C: New test. * g++.dg/gomp/map-ternary-1.C: New test. * g++.dg/gomp/member-array-2.C: New test. libgomp/ * testsuite/libgomp.c++/baseptrs-4.C: Remove commented-out cases that now work. * testsuite/libgomp.c++/baseptrs-6.C: New test. * testsuite/libgomp.c++/ind-base-1.C: New test. * testsuite/libgomp.c++/ind-base-2.C: New test. * testsuite/libgomp.c++/lvalue-tofrom-1.C: New test. * testsuite/libgomp.c++/lvalue-tofrom-2.C: New test. * testsuite/libgomp.c++/map-comma-1.C: New test. * testsuite/libgomp.c++/map-rvalue-ref-1.C: New test. * testsuite/libgomp.c++/struct-ref-1.C: New test. * testsuite/libgomp.c-c++-common/array-field-1.c: New test. * testsuite/libgomp.c-c++-common/array-of-struct-1.c: New test. * testsuite/libgomp.c-c++-common/array-of-struct-2.c: New test. |
||
|
Jakub Jelinek
|
a945c346f5 | Update copyright years. | ||
|
Julian Brown
|
144c531fe2 |
OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc
This patch has been separated out from the C++ "declare mapper" support patch. It contains just the gimplify.cc rearrangement work, mostly moving gimplification from gimplify_scan_omp_clauses to gimplify_adjust_omp_clauses for map clauses. The motivation for doing this was that we don't know if we need to instantiate mappers implicitly until the body of an offload region has been scanned, i.e. in gimplify_adjust_omp_clauses, but we also need the un-gimplified form of clauses to sort by base-pointer dependencies after mapper instantiation has taken place. The patch also reimplements the "present" clause sorting code to avoid another sorting pass on mapping nodes. This version of the patch is based on the version posted for og13, and additionally incorporates a follow-on fix for DECL_VALUE_EXPR handling in gimplify_adjust_omp_clauses: "OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622223.html Parts of: "OpenMP: OpenMP 5.2 semantics for pointers with unmapped target" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623351.html 2023-12-16 Julian Brown <julian@codesourcery.com> gcc/ * gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups. (gimplify_scan_omp_clauses): Use mapping group functionality to iterate through mapping nodes. Remove most gimplification of OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables splay tree. (gimplify_adjust_omp_clauses): Move most gimplification of OMP_CLAUSE_MAP nodes here. libgomp/ * testsuite/libgomp.fortran/target-enter-data-6.f90: Remove XFAIL. |
||
|
Julian Brown
|
d7e9ae4fa9 |
OpenMP, NVPTX: memcpy[23]D bias correction
This patch works around behaviour of the 2D and 3D memcpy operations in the CUDA driver runtime. Particularly in Fortran, the "base pointer" of an array (used for either source or destination of a host/device copy) may lie outside of data that is actually stored on the device. The fix is to make sure that we use the first element of data to be transferred instead, and adjust parameters accordingly. 2023-10-02 Julian Brown <julian@codesourcery.com> libgomp/ * plugin/plugin-nvptx.c (GOMP_OFFLOAD_memcpy2d): Adjust parameters to avoid out-of-bounds array checks in CUDA runtime. (GOMP_OFFLOAD_memcpy3d): Likewise. * testsuite/libgomp.c-c++-common/memcpyxd-bias-1.c: New test. |
||
|
Jakub Jelinek
|
000155e8ee |
libgomp: Make libgomp.c/declare-variant-1.c test x86 specific
As written earlier, this test was written with the x86 specifics in mind and adding dg-final directives for it for other arches makes it unreadable. If a declare variant call can be resolved in gimple already as in the aarch64 or gcn cases, it can be done in gcc.dg/gomp/ and I believe we have tests like that already, the point of the test is that it is not known during gimplification time which exact call should be chosen as it depends on which declare simd clone it will be in. 2023-12-18 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/declare-variant-1.c: Restrict the test to x86, drop because of that unneeded target selector from other directives and remove the aarch64 specific ones. |
||
|
Andre Vieira
|
863df360fb |
Fix tests for gomp
This is to fix testisms initially introduced by:
commit
|
||
|
Thomas Schwinge
|
bc7546e32c |
In 'libgomp.fortran/map-subarray-5.f90', restrict 'dg-output's to 'target offload_device_nonshared_as'
..., as in 'libgomp.c-c++-common/map-arrayofstruct-{2,3}.c'.
Minor fix-up for commit
|
||
|
Julian Brown
|
f5745dc142 |
OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic
This patch adds support for non-constant component offsets in "map" clauses for OpenMP (and the equivalants for OpenACC), which are not able to be sorted into order at compile time. Normally struct accesses in such clauses are gathered together and sorted into increasing address order after a "GOMP_MAP_STRUCT" node: if we have variable indices, that is no longer possible. This version of the patch scales back the previously-posted version to merely add a diagnostic for incorrect usage of component accesses with variably-indexed arrays of structs: the only permitted variant is where we have multiple indices that are the same, but we could not prove so at compile time. Rather than silently producing the wrong result for cases where the indices are in fact different, we error out (e.g., "map(dtarr(i)%arrptr, dtarr(j)%arrptr(4:8))", for different i/j). For now, multiple *constant* array indices are still supported (see map-arrayofstruct-1.c). That could perhaps be addressed with a follow-up patch, if necessary. This version of the patch renumbers the GOMP_MAP_STRUCT_UNORD kind to avoid clashing with the OpenACC "non-contiguous" dynamic array support (though that is not yet applied to mainline). 2023-08-18 Julian Brown <julian@codesourcery.com> gcc/ * gimplify.cc (extract_base_bit_offset): Add VARIABLE_OFFSET parameter. (omp_get_attachment, omp_group_last, omp_group_base, omp_directive_maps_explicitly): Add GOMP_MAP_STRUCT_UNORD support. (omp_accumulate_sibling_list): Update calls to extract_base_bit_offset. Support GOMP_MAP_STRUCT_UNORD. (omp_build_struct_sibling_lists, gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses, gimplify_omp_target_update): Add GOMP_MAP_STRUCT_UNORD support. * omp-low.cc (lower_omp_target): Add GOMP_MAP_STRUCT_UNORD support. * tree-pretty-print.cc (dump_omp_clause): Likewise. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_STRUCT_UNORD. libgomp/ * oacc-mem.c (find_group_last, goacc_enter_data_internal, goacc_exit_data_internal, GOACC_enter_exit_data): Add GOMP_MAP_STRUCT_UNORD support. * target.c (gomp_map_vars_internal): Add GOMP_MAP_STRUCT_UNORD support. Detect incorrect use of variable indexing of arrays of structs. (GOMP_target_enter_exit_data, gomp_target_task_fn): Add GOMP_MAP_STRUCT_UNORD support. * testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c: New test. * testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c: New test. * testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c: New test. * testsuite/libgomp.fortran/map-subarray-5.f90: New test. |
||
|
Julian Brown
|
7362543f00 |
OpenMP: Pointers and member mappings
This patch changes the mapping node arrangement used for array components of derived types in order to accommodate for changes made in the previous patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed derived-type members instead of "GOMP_MAP_ALWAYS_POINTER". We change the mapping nodes used for a derived-type mapping like this: type T integer, pointer, dimension(:) :: arrptr end type T type(T) :: tvar [...] !$omp target map(tofrom: tvar%arrptr) So that the nodes used look like this: 1) map(to: tvar%arrptr) --> GOMP_MAP_TO [implicit] *tvar%arrptr%data (the array data) GOMP_MAP_TO_PSET tvar%arrptr (the descriptor) GOMP_MAP_ATTACH_DETACH tvar%arrptr%data 2) map(tofrom: tvar%arrptr(3:8) --> GOMP_MAP_TOFROM *tvar%arrptr%data(3) (size 8-3+1, etc.) GOMP_MAP_TO_PSET tvar%arrptr GOMP_MAP_ATTACH_DETACH tvar%arrptr%data (bias 3, etc.) In this case, we can determine in the front-end that the whole-array/pointer mapping (1) is only needed to map the pointer -- so we drop it entirely. (Note also that we set -- early -- the OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P flag for whole-array-via-pointer mappings. See below.) In the middle end, we process mappings using the struct sibling-list handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle of the group of three mapping nodes to the proper sorted position after the GOMP_MAP_STRUCT mapping: GOMP_MAP_STRUCT tvar (len: 1) GOMP_MAP_TO_PSET tvar%arr (size: 64, etc.) <--. moved here [...] | GOMP_MAP_TOFROM *tvar%arrptr%data(3) ___| GOMP_MAP_ATTACH_DETACH tvar%arrptr%data In another case, if we have an array of derived-type values "dtarr", and mappings like: i = 1 j = 1 map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8)) We still map the same way, but this time we cannot prove that the base expressions "dtarr(i) and "dtarr(j)" are the same in the front-end. So we keep both mappings, but we move the "[implicit]" mapping of the full-array reference to the end of the clause list in gimplify.cc (by adjusting the topological sorting algorithm): GOMP_MAP_STRUCT dtvar (len: 2) GOMP_MAP_TO_PSET dtvar(i)%arrptr GOMP_MAP_TO_PSET dtvar(j)%arrptr [...] GOMP_MAP_TOFROM *dtvar(j)%arrptr%data(3) (size: 8-3+1) GOMP_MAP_ATTACH_DETACH dtvar(j)%arrptr%data GOMP_MAP_TO [implicit] *dtvar(i)%arrptr%data(1) (size: whole array) GOMP_MAP_ATTACH_DETACH dtvar(i)%arrptr%data Always moving "[implicit]" full-array mappings after array-section mappings (without that bit set) means that we'll avoid copying the whole array unnecessarily -- even in cases where we can't prove that the arrays are the same. The patch also fixes some bugs with "enter data" and "exit data" directives with this new mapping arrangement. Also now if you have mappings like this: #pragma omp target enter data map(to: dv, dv%arr(1:20)) The whole of the derived-type variable "dv" is mapped, so the GOMP_MAP_TO_PSET for the array-section mapping can be dropped: GOMP_MAP_TO dv GOMP_MAP_TO *dv%arr%data GOMP_MAP_TO_PSET dv%arr <-- deleted (array section mapping) GOMP_MAP_ATTACH_DETACH dv%arr%data To accommodate for recent changes to mapping nodes made by Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET for "exit data" directives, in favour of using the "correct" GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion. A new flag is introduced so the middle-end knows when the latter two kinds are being used specifically for an array descriptor. This version of the patch fixes "omp target exit data" handling for GOMP_MAP_DELETE, and adds pretty-printing dump output for the OMP_CLAUSE_RELEASE_DESCRIPTOR flag (for a little extra clarity). Also I noticed the handling of descriptors on *OpenACC* exit-data directives was inconsistent, so I've made those use GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the new flag in the same way as OpenMP too. In the end it doesn't actually matter to the runtime, which handles GOMP_MAP_RELEASE/GOMP_MAP_DELETE/GOMP_MAP_TO_PSET for array descriptors on OpenACC "exit data" directives the same, anyway, and doing it this way in the FE avoids needless divergence. I've added a couple of new tests (gomp/target-enter-exit-data.f90 and goacc/enter-exit-data-2.f90). 2023-12-07 Julian Brown <julian@codesourcery.com> gcc/fortran/ * dependency.cc (gfc_omp_expr_prefix_same): New function. * dependency.h (gfc_omp_expr_prefix_same): Add prototype. * gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2" union. * trans-openmp.cc (dependency.h): Include. (gfc_trans_omp_array_section): Adjust mapping node arrangement for array descriptors. Use GOMP_MAP_TO_PSET or GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DESCRIPTOR flag set. (gfc_symbol_rooted_namelist): New function. (gfc_trans_omp_clauses): Check subcomponent and subarray/element accesses elsewhere in the clause list for pointers to derived types or array descriptors, and adjust or drop mapping nodes appropriately. Adjust for changes to mapping node arrangement. (gfc_trans_oacc_executable_directive): Pass code op through. gcc/ * gimplify.cc (omp_map_clause_descriptor_p): New function. (build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base): Use above function. (omp_tsort_mapping_groups): Process nodes that have OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't. Add enter_exit_data parameter. (omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if we're mapping the whole containing derived-type variable. (omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling. Remove GOMP_MAP_ALWAYS_POINTER handling. (gimplify_scan_omp_clauses): Pass enter_exit argument to omp_tsort_mapping_groups. Don't adjust/remove GOMP_MAP_TO_PSET mappings for derived-type components here. * tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro. * tree-pretty-print.cc (dump_omp_clause): Show OMP_CLAUSE_RELEASE_DESCRIPTOR in dump output (with GOMP_MAP_TO_PSET-like syntax). gcc/testsuite/ * gfortran.dg/goacc/enter-exit-data-2.f90: New test. * gfortran.dg/goacc/finalize-1.f: Adjust scan output. * gfortran.dg/gomp/map-9.f90: Adjust scan output. * gfortran.dg/gomp/map-subarray-2.f90: New test. * gfortran.dg/gomp/map-subarray.f90: New test. * gfortran.dg/gomp/target-enter-exit-data.f90: New test. libgomp/ * testsuite/libgomp.fortran/map-subarray.f90: New test. * testsuite/libgomp.fortran/map-subarray-2.f90: New test. * testsuite/libgomp.fortran/map-subarray-3.f90: New test. * testsuite/libgomp.fortran/map-subarray-4.f90: New test. * testsuite/libgomp.fortran/map-subarray-6.f90: New test. * testsuite/libgomp.fortran/map-subarray-7.f90: New test. * testsuite/libgomp.fortran/map-subarray-8.f90: New test. * testsuite/libgomp.fortran/map-subcomponents.f90: New test. * testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for descriptor-mapping changes. Remove XFAIL. |
||
|
Julian Brown
|
5fdb150cd4 |
OpenMP/OpenACC: Rework clause expansion and nested struct handling
This patch reworks clause expansion in the C, C++ and (to a lesser
extent) Fortran front ends for OpenMP and OpenACC mapping nodes used in
GPU offloading support.
At present a single clause may be turned into several mapping nodes,
or have its mapping type changed, in several places scattered through
the front- and middle-end. The analysis relating to which particular
transformations are needed for some given expression has become quite hard
to follow. Briefly, we manipulate clause types in the following places:
1. During parsing, in c_omp_adjust_map_clauses. Depending on a set of
rules, we may change a FIRSTPRIVATE_POINTER (etc.) mapping into
ATTACH_DETACH, or mark the decl addressable.
2. In semantics.cc or c-typeck.cc, clauses are expanded in
handle_omp_array_sections (called via {c_}finish_omp_clauses, or in
finish_omp_clauses itself. The two cases are for processing array
sections (the former), or non-array sections (the latter).
3. In gimplify.cc, we build sibling lists for struct accesses, which
groups and sorts accesses along with their struct base, creating
new ALLOC/RELEASE nodes for pointers.
4. In gimplify.cc:gimplify_adjust_omp_clauses, mapping nodes may be
adjusted or created.
This patch doesn't completely disrupt this scheme, though clause
types are no longer adjusted in c_omp_adjust_map_clauses (step 1).
Clause expansion in step 2 (for C and C++) now uses a single, unified
mechanism, parts of which are also reused for analysis in step 3.
Rather than the kind-of "ad-hoc" pattern matching on addresses used to
expand clauses used at present, a new method for analysing addresses is
introduced. This does a recursive-descent tree walk on expression nodes,
and emits a vector of tokens describing each "part" of the address.
This tokenized address can then be translated directly into mapping nodes,
with the assurance that no part of the expression has been inadvertently
skipped or misinterpreted. In this way, all the variations of ways
pointers, arrays, references and component accesses might be combined
can be teased apart into easily-understood cases - and we know we've
"parsed" the whole address before we start analysis, so the right code
paths can easily be selected.
For example, a simple access "arr[idx]" might parse as:
base-decl access-indexed-array
or "mystruct->foo[x]" with a pointer "foo" component might parse as:
base-decl access-pointer component-selector access-pointer
A key observation is that support for "array" bases, e.g. accesses
whose root nodes are not structures, but describe scalars or arrays,
and also *one-level deep* structure accesses, have first-class support
in gimplify and beyond. Expressions that use deeper struct accesses
or e.g. multiple indirections were more problematic: some cases worked,
but lots of cases didn't. This patch reimplements the support for those
in gimplify.cc, again using the new "address tokenization" support.
An expression like "mystruct->foo->bar[0:10]" used in a mapping node will
translate the right-hand access directly in the front-end. The base for
the access will be "mystruct->foo". This is handled recursively in
gimplify.cc -- there may be several accesses of "mystruct"'s members
on the same directive, so the sibling-list building machinery can be
used again. (This was already being done for OpenACC, but the new
implementation differs somewhat in details, and is more robust.)
For OpenMP, in the case where the base pointer itself,
i.e. "mystruct->foo" here, is NOT mapped on the same directive, we
create a "fragile" mapping. This turns the "foo" component access
into a zero-length allocation (which is a new feature for the runtime,
so support has been added there too).
A couple of changes have been made to how mapping clauses are turned
into mapping nodes:
The first change is based on the observation that it is probably never
correct to use GOMP_MAP_ALWAYS_POINTER for component accesses (e.g. for
references), because if the containing struct is already mapped on the
target then the host version of the pointer in question will be corrupted
if the struct is copied back from the target. This patch removes all
such uses, across each of C, C++ and Fortran.
The second change is to the way that GOMP_MAP_ATTACH_DETACH nodes
are processed during sibling-list creation. For OpenMP, for pointer
components, we must map the base pointer separately from an array section
that uses the base pointer, so e.g. we must have both "map(mystruct.base)"
and "map(mystruct.base[0:10])" mappings. These create nodes such as:
GOMP_MAP_TOFROM mystruct.base
G_M_TOFROM *mystruct.base [len: 10*elemsize] G_M_ATTACH_DETACH mystruct.base
Instead of using the first of these directly when building the struct
sibling list then skipping the group using GOMP_MAP_ATTACH_DETACH,
leading to:
GOMP_MAP_STRUCT mystruct [len: 1] GOMP_MAP_TOFROM mystruct.base
we now introduce a new "mini-pass", omp_resolve_clause_dependencies, that
drops the GOMP_MAP_TOFROM for the base pointer, marks the second group
as having had a base-pointer mapping, then omp_build_struct_sibling_lists
can create:
GOMP_MAP_STRUCT mystruct [len: 1] GOMP_MAP_ALLOC mystruct.base [len: ptrsize]
This ends up working better in many cases, particularly those involving
references. (The "alloc" space is immediately overwritten by a pointer
attachment, so this is mildly more efficient than a redundant TO mapping
at runtime also.)
There is support in the address tokenizer for "arbitrary" base expressions
which aren't rooted at a decl, but that is not used as present because
such addresses are disallowed at parse time.
In the front-ends, the address tokenization machinery is mostly only
used for clause expansion and not for diagnostics at present. It could
be used for those too, which would allow more of my previous "address
inspector" implementation to be removed.
The new bits in gimplify.cc work with OpenACC also.
This version of the patch addresses several first-pass review comments
from Tobias, and fixes a few previously-missed cases for manually-managed
ragged array mappings (including cases using references). Some arbitrary
differences between handling of clause expansion for C vs. C++ have also
been fixed, and some fragments from later in the patch series have been
moved forward (where they were useful for fixing bugs). Several new
test cases have been added.
2023-11-29 Julian Brown <julian@codesourcery.com>
gcc/c-family/
* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
C_ORT_OMP_EXIT_DATA and C_ORT_ACC_TARGET.
(omp_addr_token): Add forward declaration.
(c_omp_address_inspector): New class.
* c-omp.cc (c_omp_adjust_map_clauses): Mark decls addressable here, but
do not change any mapping node types.
(c_omp_address_inspector::unconverted_ref_origin,
c_omp_address_inspector::component_access_p,
c_omp_address_inspector::check_clause,
c_omp_address_inspector::get_root_term,
c_omp_address_inspector::map_supported_p,
c_omp_address_inspector::get_origin,
c_omp_address_inspector::maybe_unconvert_ref,
c_omp_address_inspector::maybe_zero_length_array_section,
c_omp_address_inspector::expand_array_base,
c_omp_address_inspector::expand_component_selector,
c_omp_address_inspector::expand_map_clause): New methods.
(omp_expand_access_chain): New function.
gcc/c/
* c-parser.cc (c_parser_oacc_all_clauses): Add TARGET_P parameter. Use
to select region type for c_finish_omp_clauses call.
(c_parser_oacc_loop): Update calls to c_parser_oacc_all_clauses.
(c_parser_oacc_compute): Likewise.
(c_parser_omp_target_data, c_parser_omp_target_enter_data): Support
ATTACH kind.
(c_parser_omp_target_exit_data): Support DETACH kind.
(check_clauses): Handle GOMP_MAP_POINTER and GOMP_MAP_ATTACH here.
* c-typeck.cc (handle_omp_array_sections_1,
handle_omp_array_sections, c_finish_omp_clauses): Use
c_omp_address_inspector class and OMP address tokenizer to analyze and
expand map clause expressions. Fix some diagnostics. Fix "is OpenACC"
condition for C_ORT_ACC_TARGET addition.
gcc/cp/
* parser.cc (cp_parser_oacc_all_clauses): Add TARGET_P parameter. Use
to select region type for finish_omp_clauses call.
(cp_parser_omp_target_data, cp_parser_omp_target_enter_data): Support
GOMP_MAP_ATTACH kind.
(cp_parser_omp_target_exit_data): Support GOMP_MAP_DETACH kind.
(cp_parser_oacc_declare): Update call to cp_parser_oacc_all_clauses.
(cp_parser_oacc_loop): Update calls to cp_parser_oacc_all_clauses.
(cp_parser_oacc_compute): Likewise.
* pt.cc (tsubst_expr): Use C_ORT_ACC_TARGET for call to
tsubst_omp_clauses for OpenACC compute regions.
* semantics.cc (cp_omp_address_inspector): New class, derived from
c_omp_address_inspector.
(handle_omp_array_sections_1, handle_omp_array_sections,
finish_omp_clauses): Use cp_omp_address_inspector class and OMP address
tokenizer to analyze and expand OpenMP map clause expressions. Fix
some diagnostics. Support C_ORT_ACC_TARGET.
(finish_omp_target): Handle GOMP_MAP_POINTER.
gcc/fortran/
* trans-openmp.cc (gfc_trans_omp_array_section): Add OPENMP parameter.
Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for
derived type components.
(gfc_trans_omp_clauses): Update calls to gfc_trans_omp_array_section.
gcc/
* gimplify.cc (build_struct_comp_nodes): Don't process
GOMP_MAP_ATTACH_DETACH "middle" nodes here.
(omp_mapping_group): Add REPROCESS_STRUCT and FRAGILE booleans for
nested struct handling.
(omp_strip_components_and_deref, omp_strip_indirections): Remove
functions.
(omp_get_attachment): Handle GOMP_MAP_DETACH here.
(omp_group_last): Handle GOMP_MAP_*, GOMP_MAP_DETACH,
GOMP_MAP_ATTACH_DETACH groups for "exit data" of reference-to-pointer
component array sections.
(omp_gather_mapping_groups_1): Initialise reprocess_struct and fragile
fields.
(omp_group_base): Handle GOMP_MAP_ATTACH_DETACH after GOMP_MAP_STRUCT.
(omp_index_mapping_groups_1): Skip reprocess_struct groups.
(omp_get_nonfirstprivate_group, omp_directive_maps_explicitly,
omp_resolve_clause_dependencies, omp_first_chained_access_token): New
functions.
(omp_check_mapping_compatibility): Adjust accepted node combinations
for "from" clauses using release instead of alloc.
(omp_accumulate_sibling_list): Add GROUP_MAP, ADDR_TOKENS, FRAGILE_P,
REPROCESSING_STRUCT, ADDED_TAIL parameters. Use OMP address tokenizer
to analyze addresses. Reimplement nested struct handling, and
implement "fragile groups".
(omp_build_struct_sibling_lists): Adjust for changes to
omp_accumulate_sibling_list. Recalculate bias for ATTACH_DETACH nodes
after GOMP_MAP_STRUCT nodes.
(gimplify_scan_omp_clauses): Call omp_resolve_clause_dependencies. Use
OMP address tokenizer.
(gimplify_adjust_omp_clauses_1): Use build_fold_indirect_ref_loc
instead of build_simple_mem_ref_loc.
* omp-general.cc (omp-general.h, tree-pretty-print.h): Include.
(omp_addr_tokenizer): New namespace.
(omp_addr_tokenizer::omp_addr_token): New.
(omp_addr_tokenizer::omp_parse_component_selector,
omp_addr_tokenizer::omp_parse_ref,
omp_addr_tokenizer::omp_parse_pointer,
omp_addr_tokenizer::omp_parse_access_method,
omp_addr_tokenizer::omp_parse_access_methods,
omp_addr_tokenizer::omp_parse_structure_base,
omp_addr_tokenizer::omp_parse_structured_expr,
omp_addr_tokenizer::omp_parse_array_expr,
omp_addr_tokenizer::omp_access_chain_p,
omp_addr_tokenizer::omp_accessed_addr): New functions.
(omp_parse_expr, debug_omp_tokenized_addr): New functions.
* omp-general.h (omp_addr_tokenizer::access_method_kinds,
omp_addr_tokenizer::structure_base_kinds,
omp_addr_tokenizer::token_type,
omp_addr_tokenizer::omp_addr_token,
omp_addr_tokenizer::omp_access_chain_p,
omp_addr_tokenizer::omp_accessed_addr): New.
(omp_addr_token, omp_parse_expr): New.
* omp-low.cc (scan_sharing_clauses): Skip error check for references
to pointers.
* tree.h (OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED): New macro.
gcc/testsuite/
* c-c++-common/gomp/clauses-2.c: Fix error output.
* c-c++-common/gomp/target-implicit-map-2.c: Adjust scan output.
* c-c++-common/gomp/target-50.c: Adjust scan output.
* c-c++-common/gomp/target-enter-data-1.c: Adjust scan output.
* g++.dg/gomp/static-component-1.C: New test.
* gcc.dg/gomp/target-3.c: Adjust scan output.
* gfortran.dg/gomp/map-9.f90: Adjust scan output.
libgomp/
* target.c (gomp_map_pointer): Modify zero-length array section
pointer handling.
(gomp_attach_pointer): Likewise.
(gomp_map_fields_existing): Use gomp_map_0len_lookup.
(gomp_attach_pointer): Allow attaching null pointers (or Fortran
"unassociated" pointers).
(gomp_map_vars_internal): Handle zero-sized struct members. Add
diagnostic for unmapped struct pointer members.
* testsuite/libgomp.c-c++-common/baseptrs-1.c: New test.
* testsuite/libgomp.c-c++-common/baseptrs-2.c: New test.
* testsuite/libgomp.c-c++-common/baseptrs-6.c: New test.
* testsuite/libgomp.c-c++-common/baseptrs-7.c: New test.
* testsuite/libgomp.c-c++-common/ptr-attach-2.c: New test.
* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: Fix missing
"free".
* testsuite/libgomp.c-c++-common/target-implicit-map-5.c: New test.
* testsuite/libgomp.c-c++-common/target-map-zlas-1.c: New test.
* testsuite/libgomp.c++/class-array-1.C: New test.
* testsuite/libgomp.c++/baseptrs-3.C: New test.
* testsuite/libgomp.c++/baseptrs-4.C: New test.
* testsuite/libgomp.c++/baseptrs-5.C: New test.
* testsuite/libgomp.c++/baseptrs-8.C: New test.
* testsuite/libgomp.c++/baseptrs-9.C: New test.
* testsuite/libgomp.c++/ref-mapping-1.C: New test.
* testsuite/libgomp.c++/target-48.C: New test.
* testsuite/libgomp.c++/target-49.C: New test.
* testsuite/libgomp.c++/target-exit-data-reftoptr-1.C: New test.
* testsuite/libgomp.c++/target-lambda-1.C: Update for OpenMP 5.2
semantics.
* testsuite/libgomp.c++/target-this-3.C: Likewise.
* testsuite/libgomp.c++/target-this-4.C: Likewise.
* testsuite/libgomp.fortran/struct-elem-map-1.f90: Add temporary XFAIL.
* testsuite/libgomp.fortran/target-enter-data-6.f90: Likewise.
|
||
|
Andrew Stubbs
|
348874f0ba |
libgomp: basic pinned memory on Linux
Implement the OpenMP pinned memory trait on Linux hosts using the mlock syscall. Pinned allocations are performed using mmap, not malloc, to ensure that they can be unpinned safely when freed. This implementation will work OK for page-scale allocations, and finer-grained allocations will be implemented in a future patch. libgomp/ChangeLog: * allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. (MEMSPACE_VALIDATE): Add PIN. (omp_init_allocator): Use MEMSPACE_VALIDATE to check pinning. (omp_aligned_alloc): Add pinning to all MEMSPACE_* calls. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. (omp_free): Likewise. * config/linux/allocator.c: New file. * config/nvptx/allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. (MEMSPACE_VALIDATE): Add PIN. * config/gcn/allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. * libgomp.texi: Switch pinned trait to supported. (MEMSPACE_VALIDATE): Add PIN. * testsuite/libgomp.c/alloc-pinned-1.c: New test. * testsuite/libgomp.c/alloc-pinned-2.c: New test. * testsuite/libgomp.c/alloc-pinned-3.c: New test. * testsuite/libgomp.c/alloc-pinned-4.c: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> |