aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2020-08-03Removal of HSA offloading from gcc and libgompMartin Jambor26-1880/+4
This patch removes the generation of HSAIL from the compiler, the HSA offloading plugin from libgomp and the associated testsuite tests and infrastructure bits from the respective testsuites. Apart from removal of the obvious files, I removed bits that I found by searching for HSA related terms and by re-tracing my steps and looking at the patches that introduced HSA in the first place. I did not remove everything these patches brought in, for example: - the mechanism to pass offload-target specific info from the application to the offloading plugin - but the same mechanism is also used to communicate number of teams and the thread limit to all offload targets. - run_func hook in gomp_device_descr stays too, although now it is not used. If some future offload target would like the ability to refuse to offload some functions, it can use it. It is easy to remove as a follow-up if it is considered clutter, though. - configure options --with-hsa-runtime=PATH, -with-hsa-runtime-include=PATH and --with-hsa-runtime-lib=PATH rmeain because GCN uses them too. - Surprisingly, GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES (a constant from gomp-constants.h) appears in the source of the amdgcn libgomp plugin, although I tend to think that code path is not ever used and this patch certainly removes it from the compiler. Nevertheless, it seems it has potential value beyond HSAIL and so I've kept it, it can of course always be easily removed in the future of GCN folk abandon it too. - I assume constants OFFLOAD_TARGET_TYPE_HSA and GOMP_DEVICE_HSA need to stay indefinitely too just so that no future offload target picks that number. - I have kept dg-require-effective-target offload_device_nonshared_as requirement of thests which have it. It is quite probable I missed some small HSA artifacts but those should be easy to remove later as we find them. include/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * gomp-constants.h (GOMP_VERSION_HSA): Remove. gcc/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * hsa-brig-format.h: Moved to brig/brigfrontend. * hsa-brig.c: Removed. * hsa-builtins.def: Likewise. * hsa-common.c: Likewise. * hsa-common.h: Likewise. * hsa-dump.c: Likewise. * hsa-gen.c: Likewise. * hsa-regalloc.c: Likewise. * ipa-hsa.c: Likewise. * omp-grid.c: Likewise. * omp-grid.h: Likewise. * Makefile.in (BUILTINS_DEF): Remove hsa-builtins.def. (OBJS): Remove hsa-common.o, hsa-gen.o, hsa-regalloc.o, hsa-brig.o, hsa-dump.o, ipa-hsa.c and omp-grid.o. (GTFILES): Removed hsa-common.c and omp-expand.c. * builtins.def: Remove processing of hsa-builtins.def. (DEF_HSA_BUILTIN): Remove. * common.opt (flag_disable_hsa): Remove. (-Whsa): Ignore. * config.in (ENABLE_HSA): Removed. * configure.ac: Removed handling configuration for hsa offloading. (ENABLE_HSA): Removed. * configure: Regenerated. * doc/install.texi (--enable-offload-targets): Remove hsa from the example. (--with-hsa-runtime): Reword to reference any HSA run-time, not specifically HSA offloading. * doc/invoke.texi (Option Summary): Remove -Whsa. (Warning Options): Likewise. (Optimize Options): Remove hsa-gen-debug-stores. * doc/passes.texi (Regular IPA passes): Remove section on IPA HSA pass. * gimple-low.c (lower_stmt): Remove GIMPLE_OMP_GRID_BODY case. * gimple-pretty-print.c (dump_gimple_omp_for): Likewise. (dump_gimple_omp_block): Likewise. (pp_gimple_stmt_1): Likewise. * gimple-walk.c (walk_gimple_stmt): Likewise. * gimple.c (gimple_build_omp_grid_body): Removed function. (gimple_copy): Remove GIMPLE_OMP_GRID_BODY case. * gimple.def (GIMPLE_OMP_GRID_BODY): Removed. * gimple.h (gf_mask): Removed GF_OMP_PARALLEL_GRID_PHONY, OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY, GF_OMP_FOR_GRID_INTRA_GROUP, GF_OMP_FOR_GRID_GROUP_ITER and GF_OMP_TEAMS_GRID_PHONY. Renumbered GF_OMP_FOR_KIND_SIMD and GF_OMP_TEAMS_HOST. (gimple_build_omp_grid_body): Removed declaration. (gimple_has_substatements): Remove GIMPLE_OMP_GRID_BODY case. (gimple_omp_for_grid_phony): Removed. (gimple_omp_for_set_grid_phony): Likewise. (gimple_omp_for_grid_intra_group): Likewise. (gimple_omp_for_grid_intra_group): Likewise. (gimple_omp_for_grid_group_iter): Likewise. (gimple_omp_for_set_grid_group_iter): Likewise. (gimple_omp_parallel_grid_phony): Likewise. (gimple_omp_parallel_set_grid_phony): Likewise. (gimple_omp_teams_grid_phony): Likewise. (gimple_omp_teams_set_grid_phony): Likewise. (CASE_GIMPLE_OMP): Remove GIMPLE_OMP_GRID_BODY case. * lto-section-in.c (lto_section_name): Removed hsa. * lto-streamer.h (lto_section_type): Removed LTO_section_ipa_hsa. * lto-wrapper.c (compile_images_for_offload_targets): Remove special handling of hsa. * omp-expand.c: Do not include hsa-common.h and gt-omp-expand.h. (parallel_needs_hsa_kernel_p): Removed. (grid_launch_attributes_trees): Likewise. (grid_launch_attributes_trees): Likewise. (grid_create_kernel_launch_attr_types): Likewise. (grid_insert_store_range_dim): Likewise. (grid_get_kernel_launch_attributes): Likewise. (get_target_arguments): Remove code passing HSA grid sizes. (grid_expand_omp_for_loop): Remove. (grid_arg_decl_map): Likewise. (grid_remap_kernel_arg_accesses): Likewise. (grid_expand_target_grid_body): Likewise. (expand_omp): Remove call to grid_expand_target_grid_body. (omp_make_gimple_edges): Remove GIMPLE_OMP_GRID_BODY case. * omp-general.c: Do not include hsa-common.h. (omp_maybe_offloaded): Do not check for HSA offloading. (omp_context_selector_matches): Likewise. * omp-low.c: Do not include hsa-common.h and omp-grid.h. (build_outer_var_ref): Remove handling of GIMPLE_OMP_GRID_BODY. (scan_sharing_clauses): Remove handling of OMP_CLAUSE__GRIDDIM_. (scan_omp_parallel): Remove handling of the phoney variant. (check_omp_nesting_restrictions): Remove handling of GIMPLE_OMP_GRID_BODY and GF_OMP_FOR_KIND_GRID_LOOP. (scan_omp_1_stmt): Remove handling of GIMPLE_OMP_GRID_BODY. (lower_omp_for_lastprivate): Remove handling of gridified loops. (lower_omp_for): Remove phony loop handling. (lower_omp_taskreg): Remove phony construct handling. (lower_omp_teams): Likewise. (lower_omp_grid_body): Removed. (lower_omp_1): Remove GIMPLE_OMP_GRID_BODY case. (execute_lower_omp): Do not call omp_grid_gridify_all_targets. * opts.c (common_handle_option): Do not handle hsa when processing OPT_foffload_. * params.opt (hsa-gen-debug-stores): Remove. * passes.def: Remove pass_ipa_hsa and pass_gen_hsail. * timevar.def: Remove TV_IPA_HSA. * toplev.c: Do not include hsa-common.h. (compile_file): Do not call hsa_output_brig. * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE__GRIDDIM_. (tree_omp_clause): Remove union field dimension. * tree-nested.c (convert_nonlocal_omp_clauses): Remove the OMP_CLAUSE__GRIDDIM_ case. (convert_local_omp_clauses): Likewise. * tree-pass.h (make_pass_gen_hsail): Remove declaration. (make_pass_ipa_hsa): Likewise. * tree-pretty-print.c (dump_omp_clause): Remove GIMPLE_OMP_GRID_BODY case. * tree.c (omp_clause_num_ops): Remove the element corresponding to OMP_CLAUSE__GRIDDIM_. (omp_clause_code_name): Likewise. (walk_tree_1): Remove GIMPLE_OMP_GRID_BODY case. * tree.h (OMP_CLAUSE__GRIDDIM__DIMENSION): Remove. (OMP_CLAUSE__GRIDDIM__SIZE): Likewise. (OMP_CLAUSE__GRIDDIM__GROUP): Likewise. gcc/fortran/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * f95-lang.c (gfc_init_builtin_functions): Remove processing of hsa-builtins.def. gcc/brig/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * brigfrontend/brig-util.h (hsa_type_packed_p): Declared. * brigfrontend/brig-util.cc (hsa_type_packed_p): Moved here from removed gcc/hsa-common.c. libgomp/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * plugin/Makefrag.am: Remove configuration of HSA plugin. * aclocal.m4: Regenerated. * Makefile.in: Regenerated. * config.h.in: Regenerated. * configure: Regenerated. * plugin/configfrag.ac: Likewise. * plugin/hsa_ext_finalize.h: Removed. * plugin/plugin-hsa.c: Likewise. * testsuite/Makefile.in: Regenerated. * testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type): Remove hsa case. (check_effective_target_hsa_offloading_selected_nocache): Removed (check_effective_target_hsa_offloading_selected): Likewise. (libgomp_init): Do not add -Wno-hsa to additional_flags. * testsuite/libgomp.hsa.c/alloca-1.c: Removed test. * testsuite/libgomp.hsa.c/bitfield-1.c: Likewise. * testsuite/libgomp.hsa.c/bits-insns.c: Likewise. * testsuite/libgomp.hsa.c/builtins-1.c: Likewise. * testsuite/libgomp.hsa.c/c.exp: Likewise. * testsuite/libgomp.hsa.c/complex-1.c: Likewise. * testsuite/libgomp.hsa.c/complex-align-2.c: Likewise. * testsuite/libgomp.hsa.c/formal-actual-args-1.c: Likewise. * testsuite/libgomp.hsa.c/function-call-1.c: Likewise. * testsuite/libgomp.hsa.c/get-level-1.c: Likewise. * testsuite/libgomp.hsa.c/gridify-1.c: Likewise. * testsuite/libgomp.hsa.c/gridify-2.c: Likewise. * testsuite/libgomp.hsa.c/gridify-3.c: Likewise. * testsuite/libgomp.hsa.c/gridify-4.c: Likewise. * testsuite/libgomp.hsa.c/memory-operations-1.c: Likewise. * testsuite/libgomp.hsa.c/pr69568.c: Likewise. * testsuite/libgomp.hsa.c/pr82416.c: Likewise. * testsuite/libgomp.hsa.c/rotate-1.c: Likewise. * testsuite/libgomp.hsa.c/staticvar.c: Likewise. * testsuite/libgomp.hsa.c/switch-1.c: Likewise. * testsuite/libgomp.hsa.c/switch-branch-1.c: Likewise. * testsuite/libgomp.hsa.c/switch-sbr-2.c: Likewise. * testsuite/libgomp.hsa.c/tiling-1.c: Likewise. * testsuite/libgomp.hsa.c/tiling-2.c: Likewise. gcc/testsuite/ChangeLog: 2020-07-24 Martin Jambor <mjambor@suse.cz> * lib/target-supports.exp (check_effective_target_offload_hsa): Removed. * c-c++-common/gomp/gridify-1.c: Removed test. * c-c++-common/gomp/gridify-2.c: Likewise. * c-c++-common/gomp/gridify-3.c: Likewise. * c-c++-common/gomp/hsa-indirect-call-1.c: Likewise. * gfortran.dg/gomp/gridify-1.f90: Likewise. * gcc.dg/gomp/gomp.exp: Do not pass -Wno-hsa to tests. * g++.dg/gomp/gomp.exp: Likewise. * gfortran.dg/gomp/gomp.exp: Likewise.
2020-07-27openacc: Deep copy attach/detach should not affect reference countsJulian Brown5-2/+284
Attach and detach operations are not supposed to affect structural or dynamic reference counts for OpenACC. Previously they did so, which led to subtle problems in some circumstances. We can avoid reference-counting attach/detach operations by extending and slightly repurposing the do_detach field in target_var_desc. It is now called is_attach to better reflect its new role. 2020-07-27 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> libgomp/ * libgomp.h (struct target_var_desc): Rename do_detach field to is_attach. * oacc-mem.c (goacc_exit_datum_1): Add assert. Don't set finalize for GOMP_MAP_FORCE_DETACH. Update checking to use is_attach field. (goacc_enter_data_internal): Don't affect reference counts for attach mappings. (goacc_exit_data_internal): Don't affect reference counts for detach mappings. * target.c (gomp_map_vars_existing): Don't affect reference counts for attach mappings. (gomp_map_vars_internal): Set renamed is_attach flag unconditionally to mark attach mappings. (gomp_unmap_vars_internal): Use is_attach flag to prevent affecting reference count for attach mappings. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Mark test as shouldfail. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust to fail gracefully in no-finalize mode. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-24[testsuite] Unset 'offload_target' after useThomas Schwinge3-0/+3
..., so that we don't leak this into '*.exp' files running later. This is relevant after commit efc16503ca10bc0e934e0bace5777500e4dc757a "handle dumpbase in offloading, adjust testsuite" -- I was confused why in a (simplified) testing sequence as follows: default 'libgomp.c/c.exp' default 'libgomp.oacc-c/c.exp' '-m32' 'libgomp.c/c.exp' '-m32' 'libgomp.oacc-c/c.exp' ..., the "'-m32' 'libgomp.c/c.exp'" variant would not execute any offloading dump scanning. The reason is that the "default 'libgomp.oacc-c/c.exp'" variant ends with 'offload_target=disable' set, so that's what the "'-m32' 'libgomp.c/c.exp'" variant would then see, in particular 'gcc/testsuite/lib/scanoffload.exp:scoff'. libgomp/ * testsuite/libgomp.oacc-c++/c++.exp: Unset 'offload_target' after use. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
2020-07-23openacc: Remove unnecessary detach finalizationJulian Brown1-0/+28
The call to gomp_detach_pointer in gomp_unmap_vars_internal does not need to force finalization, and doing so may mask mismatched pointer attachments/detachments. This patch removes the forcing. 2020-07-16 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> libgomp/ * target.c (gomp_unmap_vars_internal): Remove unnecessary forcing of finalization for detach operation. * testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-23libomp: Add omp_depend_kind to omp_lib.{f90,h}Tobias Burnus1-0/+1
gcc/fortran/ChangeLog: * intrinsic.texi (OMP_LIB_KINDS): Add omp_depend_kind. libgomp/ChangeLog: * configure.ac: Add OMP_DEPEND_KIND and OMP_INT128_SIZE. * libgomp_f.h.in (omp_check_defines): Check whether sizeof of determined Fortran kind and C typedef match. * omp_lib.f90.in: Add omp_depened_kind. * omp_lib.h.in: Likewise; fix omp_alloctrait_key_kind. * configure: Regenerate. * Makefile.in: Regenerate. * testsuite/Makefile.in: Regenerate.
2020-07-22critical-hint-*.{c,f90}: Move from gcc/testsuite to libgomp/testsuiteTobias Burnus4-0/+248
libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/critical-hint-1.c: New; moved from gcc/testsuite/c-c++-common/gomp/. * testsuite/libgomp.c-c++-common/critical-hint-2.c: Likewise. * testsuite/libgomp.fortran/critical-hint-1.f90: New; moved from gcc/testsuite/gfortran.dg/gomp/. * testsuite/libgomp.fortran/critical-hint-2.f90: Likewise. gcc/testsuite/ChangeLog: * c-c++-common/gomp/critical-hint-1.c: Moved to libgomp/. * c-c++-common/gomp/critical-hint-2.c: Moved to libgomp/. * gfortran.dg/gomp/critical-hint-1.f90: Moved to libgomp/. * gfortran.dg/gomp/critical-hint-2.f90: Moved to libgomp/.
2020-07-18x86-64: Define ASM_OUTPUT_ALIGNED_DECL_LOCALH.J. Lu1-0/+18
Define ASM_OUTPUT_ALIGNED_DECL_LOCAL for large local common symbol. gcc/ PR target/95620 * config/i386/x86-64.h (ASM_OUTPUT_ALIGNED_DECL_LOCAL): New. libgomp/ PR target/95620 * testsuite/libgomp.c/pr95620.c: New test.
2020-07-16openacc: Fix standalone attach for Fortran assumed-shape array pointersJulian Brown2-0/+71
This patch makes it so that an "attach" operation for a Fortran pointer with an array descriptor copies that array descriptor to the target, and similarly that detach operations release the array descriptor. 2020-07-16 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Rework OpenACC attach/detach handling for arrays with descriptors. gcc/testsuite/ * gfortran.dg/goacc/attach-descriptor.f90: New test. libgomp/ * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: New test. * testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-15libgomp.fortran/alloc-1.F90: Fix testcase for 32bit size_tTobias Burnus1-7/+12
libgomp/ChangeLog: * testsuite/libgomp.fortran/alloc-1.F90: Use c_size_t to avoid conversion on 32bit systems from 32bit to 64bit due to -fdefault-integer-8.
2020-07-15libgomp.fortran/struct-elem-map-1.f90: Add char kind=4 testsTobias Burnus1-40/+120
As the Fortran PR 95837 has been fixed, the test could be be added. libgomp/ChangeLog: * testsuite/libgomp.fortran/struct-elem-map-1.f90: Remove unused variables; add character(kind=4) tests; update TODO comment.
2020-07-15libgomp: Add Fortran routine support for allocatorsTobias Burnus7-2/+273
libgomp/ChangeLog: * allocator.c: Add ialias for omp_init_allocator and omp_destroy_allocator. * configure.ac: Set INTPTR_T_KIND. * configure: Regenerate. * Makefile.in: Regenerate. * testsuite/Makefile.in: Regenerate. * fortran.c (omp_init_allocator_, omp_destroy_allocator_, omp_set_default_allocator_, omp_get_default_allocator_): New functions and ialias_redirect. * icv.c: Add ialias for omp_set_default_allocator and omp_get_default_allocator. * libgomp.map (OMP_5.0.1): Add omp_init_allocator_, omp_destroy_allocator_, omp_set_default_allocator_ and omp_get_default_allocator_. * omp_lib.f90.in: Add allocator traits parameters, declare allocator routines and add related kind parameters. * omp_lib.h.in: Likewise. * testsuite/libgomp.c-c++-common/alloc-2.c: Fix sizeof. * testsuite/libgomp.fortran/alloc-1.F90: New test. * testsuite/libgomp.fortran/alloc-2.F90: New test. * testsuite/libgomp.fortran/alloc-3.F: New test. * testsuite/libgomp.fortran/alloc-4.f90: New test. * testsuite/libgomp.fortran/alloc-5.f90: New test.
2020-07-14libgomp: Fix hang when profiling OpenACC programs with CUDA 9.0 nvprofKwok Cheung Yeung1-0/+80
The version of nvprof in CUDA 9.0 causes a hang when used to profile an OpenACC program. This is because it calls acc_get_device_type from a callback called during device initialization, which then attempts to acquire acc_device_lock while it is already taken, resulting in deadlock. This works around the issue by returning acc_device_none from acc_get_device_type without attempting to acquire the lock when initialization has not completed yet. 2020-07-14 Tom de Vries <tom@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * oacc-init.c (acc_init_state_lock, acc_init_state, acc_init_thread): New variable. (acc_init_1): Set acc_init_thread to pthread_self (). Set acc_init_state to initializing at the start, and to initialized at the end. (self_initializing_p): New function. (acc_get_device_type): Return acc_device_none if called by thread that is currently executing acc_init_1. * libgomp.texi (acc_get_device_type): Update documentation. (Implementation Status and Implementation-Defined Behavior): Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-2.c: New.
2020-07-14[OpenMP, Fortran] Add structure/derived-type element mappingTobias Burnus1-0/+331
gcc/fortran/ChangeLog: * openmp.c (gfc_match_omp_clauses): Match also derived-type component refs in OMP_CLAUSE_MAP. (resolve_omp_clauses): Resolve those. * trans-openmp.c (gfc_trans_omp_array_section, gfc_trans_omp_clauses): Handle OpenMP structure-element mapping. (gfc_trans_oacc_construct, gfc_trans_oacc_executable_directive, (gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Update add openacc=true in gfc_trans_omp_clauses call. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/finalize-1.f: Update dump scan pattern. * gfortran.dg/gomp/map-1.f90: Update dg-error. * gfortran.dg/gomp/map-2.f90: New test. libgomp/ChangeLog: * testsuite/libgomp.fortran/struct-elem-map-1.f90: New test.
2020-07-14[Fortran, OpenMP] Fix allocatable-components check (PR67311)Tobias Burnus1-0/+41
gcc/fortran/ChangeLog: PR fortran/67311 * trans-openmp.c (gfc_has_alloc_comps): Return false also for pointers to arrays. libgomp/ChangeLog: PR fortran/67311 * testsuite/libgomp.fortran/target-map-1.f90: New test.
2020-07-14openmp: Adjust outer bounds of non-rect loopsJakub Jelinek1-0/+230
In loops like: #pragma omp parallel for collapse(2) for (i = -4; i < 8; i++) for (j = 3 * i; j > 2 * i; j--) for some outer loop iterations there are no inner loop iterations at all, the condition is false. In order to use Summæ Potestate to count number of iterations or to transform the logical iteration number to actual iterator values using quadratic non-equation root discovery the outer iterator range needs to be adjusted, such that the inner loop has at least one iteration for each of the outer loop iterator value in the reduced range. Sometimes this adjustment is done at the start of the range, at other times at the end. This patch implements it during the compile time number of loop computation (if all expressions are compile time constants). 2020-07-14 Jakub Jelinek <jakub@redhat.com> * omp-general.h (struct omp_for_data): Add adjn1 member. * omp-general.c (omp_extract_for_data): For non-rect loop, punt on count computing if n1, n2 or step are not INTEGER_CST earlier. Narrow the outer iterator range if needed so that non-rect loop has at least one iteration for each outer range iteration. Compute adjn1. * omp-expand.c (expand_omp_for_init_vars): Use adjn1 if non-NULL instead of the outer loop's n1. * testsuite/libgomp.c/loop-21.c: New test.
2020-07-13openacc: Don't strip TO_PSET/POINTER for enter/exit dataJulian Brown1-0/+97
OpenACC 2.6 specifies that the array descriptor (when present) must be copied to the target before attaching pointers in Fortran. This patch reverses the stripping of GOMP_MAP_TO_PSET and GOMP_MAP_POINTER that was introduced by the "OpenACC reference count overhaul" patch. 2020-07-10 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/ * gimplify.c (gimplify_scan_omp_clauses): Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data directives (see also PR92929). gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update expected dump output. libgomp/ * testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-10openacc: Adjust dynamic reference count semanticsJulian Brown14-63/+150
This patch adjusts how dynamic reference counts work so that they match the semantics of the source program more closely, instead of representing "excess" reference counts beyond those that represent pointers in the internal libgomp splay-tree data structure. This allows some corner cases to be handled more gracefully. 2020-07-10 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> libgomp/ * libgomp.h (struct splay_tree_key_s): Change virtual_refcount to dynamic_refcount. (struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA. * oacc-mem.c (acc_map_data): Substitute virtual_refcount for dynamic_refcount. (acc_unmap_data): Update comment. (goacc_map_var_existing, goacc_enter_datum): Adjust for dynamic_refcount semantics. (goacc_exit_datum_1, goacc_exit_datum): Re-add some error checking. Adjust for dynamic_refcount semantics. (goacc_enter_data_internal): Implement "present" case of dynamic memory-map handling here. Update "non-present" case for dynamic_refcount semantics. (goacc_exit_data_internal): Use goacc_exit_datum_1. * target.c (gomp_map_vars_internal): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA handling. Update for dynamic_refcount handling. (gomp_unmap_vars_internal): Remove virtual_refcount handling. (gomp_load_image_to_device): Substitute dynamic_refcount for virtual_refcount. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAILs. * testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs and trace output. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Remove trace output. * testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90: New test. * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c: Remove stale comment. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Adjust XFAIL. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-09openacc: Set bias to zero for explicit attach/detach clauses in C and C++Julian Brown2-0/+94
This is a fix for the pointer (or array) size inadvertently being used for the bias with attach and detach mapping kinds, for both C and C++. 2020-07-09 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/c/ PR middle-end/95270 * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero for standalone attach/detach clauses. gcc/cp/ PR middle-end/95270 * semantics.c (finish_omp_clauses): Likewise. include/ PR middle-end/95270 * gomp-constants.h (gomp_map_kind): Expand comment for attach/detach mapping kinds. gcc/testsuite/ PR middle-end/95270 * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero bias. libgomp/ PR middle-end/95270 * testsuite/libgomp.oacc-c-c++-common/pr95270-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr95270-2.c: New test.
2020-07-09openmp: Optimize triangular loop logical iterator to actual iterators ↵Jakub Jelinek2-0/+170
computation using search for quadratic equation root(s) This patch implements the optimized logical to actual iterators computation for triangular loops. I have a rough implementation using integers, but this one uses floating point. There is a small problem that -fopenmp programs aren't linked with -lm, so it does it only if the hw has sqrt optab (and uses ifn rather than __builtin_sqrt because it obviously doesn't need errno handling etc.). Do you think it is ok this way, or should I use the integral computation using inlined isqrt (we have inequation of the form start >= x * t10 + t11 * (((x - 1) * x) / 2) where t10 and t11 are signed long long values and start unsigned long long, and the division by 2 actually is a problem for accuracy in some cases, so if we do it in integral, we need to do actually long long t12 = 2 * t10 - t11; unsigned long long t13 = t12 * t12 + start * 8 * t11; unsigned long long isqrt_ = isqrtull (t13); long long x = (((long long) isqrt_ - t12) / t11) >> 1; with careful overflow checking on all the computations before isqrtull (and on overflows use the fallback implementation). 2020-07-09 Jakub Jelinek <jakub@redhat.com> * omp-general.h (struct omp_for_data): Add min_inner_iterations and factor members. * omp-general.c (omp_extract_for_data): Initialize them and remember them in OMP_CLAUSE_COLLAPSE_COUNT if needed and restore from there. * omp-expand.c (expand_omp_for_init_counts): Fix up computation of counts[fd->last_nonrect] if fd->loop.n2 is INTEGER_CST. (expand_omp_for_init_vars): For fd->first_nonrect + 1 == fd->last_nonrect loops with for now INTEGER_CST fd->loop.n2 find quadratic equation roots instead of using fallback method when possible. * testsuite/libgomp.c/loop-19.c: New test. * testsuite/libgomp.c/loop-20.c: New test.
2020-07-03[OpenACC] Revert always-copyfrom behavior for 'GOMP_MAP_FORCE_FROM' in ↵Thomas Schwinge1-1/+9
'libgomp/oacc-mem.c:goacc_exit_data_internal' As done for 'GOMP_MAP_FROM', also for 'GOMP_MAP_FORCE_FROM' we should only 'gomp_copy_dev2host' if 'n->refcount == 0'. This had gotten altered in commit 378da98fcc907d05002bcd3d6ff7951f0cf485e5 (r279621) "OpenACC reference count overhaul". libgomp/ * oacc-mem.c (goacc_exit_data_internal): Revert always-copyfrom behavior for 'GOMP_MAP_FORCE_FROM'. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Adjust XFAIL.
2020-06-30[testsuite] Replace fragile 'scan-assembler' with 'scan-offload-rtl' in ↵Thomas Schwinge5-10/+10
'libgomp.oacc-c-c++-common/pr85381*.c' These test cases use directives similar to: /* { dg-additional-options "-save-temps" } */ /* { dg-final { scan-assembler-times "bar.sync" 2 } } */ This expects to scan the PTX offloading compilation assembler code (not host code!), expecting that nvptx offloading code assembly is produced after the host code, and thus overwrites the latter file. (Yes, that's certainly ugly/fragile...) ..., and this broke with recent commit 1dedc12d186a110854537e1279b4e6c29f2df35a "revamp dump and aux output names" plus fix-up commit commit efc16503ca10bc0e934e0bace5777500e4dc757a "handle dumpbase in offloading, adjust testsuite" (short summary: file names changed), so let's finally make that robust. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Replace fragile 'scan-assembler' with 'scan-offload-rtl'. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise.
2020-06-27openmp: Non-rectangular loop support for non-composite worksharing loops and ↵Jakub Jelinek2-0/+434
distribute This implements the fallback mentioned in https://gcc.gnu.org/pipermail/gcc/2020-June/232874.html Special cases for triangular loops etc. to follow later, also composite constructs not supported yet (need to check the passing of temporaries around) and lastprivate might not give the same answers as serial loop if the last innermost body iteration isn't the last one for some of the outer loops (that will need to be solved separately together with rectangular loops that have no innermost body iterations, but some of the outer loops actually iterate). Also, simd needs work. 2020-06-27 Jakub Jelinek <jakub@redhat.com> * omp-general.h (struct omp_for_data_loop): Add non_rect_referenced member, move outer member. (struct omp_for_data): Add first_nonrect and last_nonrect members. * omp-general.c (omp_extract_for_data): Initialize first_nonrect, last_nonrect and non_rect_referenced members. * omp-expand.c (expand_omp_for_init_counts): Handle non-rectangular loops. (expand_omp_for_init_vars): Add nonrect_bounds parameter. Handle non-rectangular loops. (extract_omp_for_update_vars): Likewise. (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk, expand_omp_simd, expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): Adjust expand_omp_for_init_vars and extract_omp_for_update_vars callers. (expand_omp_for): Don't sorry on non-composite worksharing-loop or distribute. * testsuite/libgomp.c/loop-17.c: New test. * testsuite/libgomp.c/loop-18.c: New test.
2020-06-26c++: Change the default dialect to C++17.Marek Polacek1-1/+2
Since GCC 9, C++17 support is no longer experimental. It was too late to change the default C++ dialect to C++17 in GCC 10, but I think now it's time to pull the trigger (C++14 was made the default in GCC 6.1). We're still missing two C++17 library features, but that shouldn't stop us. See <https://gcc.gnu.org/onlinedocs/libstdc++/manual/status.html#status.iso.2017> and <https://gcc.gnu.org/projects/cxx-status.html#cxx17> for the C++17 status. I won't list all C++17 features here, but just a few heads-up: - trigraphs were removed (hardly anyone cares, unless your keyboard is missing the # key), - operator++(bool) was removed (so some tests now run in C++14 and down only), - the keyword register was removed (some legacy code might trip on this), - noexcept specification is now part of the type system and C++17 does not allow dynamic exception specifications anymore (the empty throw specification is still available, but it is deprecated), - the evaluation order rules are different in C++17, - static constexpr data members are now implicitly inline (which makes them definitions), - C++17 requires guaranteed copy elision, meaning that a copy/move constructor call might be elided completely. That means that if something relied on a constructor being instantiated via e.g. copying a function parameter, it might now fail. I'll post an update for cxx-status.html and add a new caveat to changes.html once this is in. gcc/ChangeLog: * doc/invoke.texi (C Dialect Options): Adjust -std default for C++. * doc/standards.texi (C Language): Correct the default dialect. (C++ Language): Update the default for C++ to gnu++17. gcc/c-family/ChangeLog: * c-opts.c (c_common_init_options): Default to gnu++17. gcc/testsuite/ChangeLog: * c-c++-common/torture/vector-subscript-3.c: In C++17, define away the keyword register. * g++.dg/cpp1z/attributes-enum-1a.C: Only run pre-C++17. * g++.dg/cpp1z/fold7a.C: Likewise. * g++.dg/cpp1z/nontype3a.C: Likewise. * g++.dg/cpp1z/utf8-2a.C: Likewise. * g++.dg/parse/error11.C: Update expected diagnostics for C++17. * g++.dg/torture/pr34850.C: Add -Wno-attribute-warning. * g++.dg/torture/pr49394.C: In C++17, use noexcept(false). * g++.dg/torture/pr82154.C: Use -std=c++14. * lib/target-supports.exp: Set to C++17. * obj-c++.dg/try-catch-9.mm: Use -Wno-register. libgomp/ChangeLog: * testsuite/libgomp.c++/atomic-3.C: Use -std=gnu++14.
2020-06-23handle dumpbase in offloading, adjust testsuiteAlexandre Oliva2-8/+1
Pass dumpbase on to mkoffloads and their offload-target compiler runs, using different suffixes for different offloading targets. Obey -save-temps in naming temporary files while at that. Adjust the testsuite offload dump scanning machinery to look for dump files named under the new conventions, iterating internally over all configured offload targets, or recognizing libgomp's testsuite's own iteration. for gcc/ChangeLog * collect-utils.h (dumppfx): New. * collect-utils.c (dumppfx): Likewise. * lto-wrapper.c (run_gcc): Set global dumppfx. (compile_offload_image): Pass a -dumpbase on to mkoffload. * config/nvptx/mkoffload.c (ptx_dumpbase): New. (main): Handle incoming -dumpbase. Set ptx_dumpbase. Obey save_temps. (compile_native): Pass -dumpbase et al to compiler. * config/gcn/mkoffload.c (gcn_dumpbase): New. (main): Handle incoming -dumpbase. Set gcn_dumpbase. Obey save_temps. Pass -dumpbase et al to offload target compiler. (compile_native): Pass -dumpbase et al to compiler. for gcc/testsuite/ChangeLog * lib/scanoffload.exp: New. * lib/scanoffloadrtl.exp: Load it. Replace ".o" with "" globally, and use scanoffload's scoff wrapper to fill it in. * lib/scanoffloadtree.exp: Likewise. for libgomp/ChangeLog * testsuite/lib/libgomp.exp: Load gcc lib scanoffload.exp. * testsuite/lib/libgomp-dg.exp: Drop now-obsolete -save-temps.
2020-06-18Add 'dg-do run' to 'libgomp.fortran/use_device_ptr-optional-3.f90' [PR94848]Thomas Schwinge1-0/+1
Fix-up for r279858/commit f760c0c77fe350616da9dbeaea16442b0acfb09c "Fortran] OpenMP/OpenACC – fix more issues with OPTIONAL". With offloading enabled, we then saw: PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O0 (test for excess errors) PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O0 execution test PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O1 (test for excess errors) PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O1 execution test FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O2 (test for excess errors) UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O2 compilation failed to produce executable FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (test for excess errors) UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions compilation failed to produce executable FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -g (test for excess errors) UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -g compilation failed to produce executable FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -Os (test for excess errors) UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -Os compilation failed to produce executable ... due to: /tmp/cciVc43I.o:(.gnu.offload_vars+0x10): undefined reference to `A.12.4064' [...] ..., but after the recent PR94848, PR95551 changes, that problem is now gone. libgomp/ PR lto/94848 * testsuite/libgomp.fortran/use_device_ptr-optional-3.f90: Add 'dg-do run'.
2020-06-16OpenACC/Fortran: permit 'routine' inside PURETobias Burnus1-0/+52
gcc/fortran/ChangeLog * parse.c (decode_oacc_directive): Permit 'acc routine' also inside pure procedures. * openmp.c (gfc_match_oacc_routine): Inside pure procedures do not permit gang, worker or vector clauses. libgomp/ChangeLog: * testsuite/libgomp.oacc-fortran/routine-10.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/pure-elemental-procedures-2.f90: New test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-06-08openmp: ensure variables in offload table are streamed out (PRs 94848 + 95551)Tobias Burnus1-0/+32
gcc/ChangeLog: PR lto/94848 PR middle-end/95551 * omp-offload.c (add_decls_addresses_to_decl_constructor, omp_finish_file): Skip removed items. * lto-cgraph.c (output_offload_tables): Likewise; set force_output to this node for variables and functions. libgomp/ChangeLog: PR lto/94848 PR middle-end/95551 * testsuite/libgomp.fortran/target-var.f90: New test.
2020-06-05Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'Julian Brown2-0/+82
libgomp/ * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-06-04[OpenACC 'exit data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' ↵Thomas Schwinge1-26/+67
entries Currently, we don't at all evaluate 'copyfrom' for 'GOMP_MAP_STRUCT' entries. Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code. libgomp/ * oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Evaluate 'copyfrom' individually for each entry. * testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update.
2020-06-04[OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' ↵Thomas Schwinge2-47/+146
entries Currently, we don't at all evaluate 'finalize' for 'GOMP_MAP_STRUCT' entries. Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code. libgomp/ * oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Evaluate 'finalize' individually for each entry. * testsuite/libgomp.oacc-c-c++-common/struct-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: Remove file.
2020-06-04Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7,8}.c'Thomas Schwinge2-4/+4
libgomp/ * testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: Fix 'sizeof' usage. * testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: Likewise.
2020-06-04[OpenACC] Repair/restore 'is_tgt_unmapped' checkingThomas Schwinge7-26/+79
libgomp/ * oacc-mem.c (goacc_exit_datum): Repair 'is_tgt_unmapped' checking. (acc_unmap_data, goacc_exit_data_internal): Restore 'is_tgt_unmapped' checking. * testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: New file. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise. Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-04Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854]Thomas Schwinge1-17/+47
libgomp/ PR libgomp/92854 * testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: Extend some more.
2020-06-04[OpenACC] XFAIL behavior of over-eager 'finalize' clauseThomas Schwinge10-0/+336
libgomp/ * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: XFAIL behavior of over-eager 'finalize' clause. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: New file. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90: Likewise.
2020-06-04'libgomp.oacc-fortran/{error_,}stop-{1,2,3}.f': initialize before the checkpointThomas Schwinge6-0/+24
If, for example, GCC is configured such that 'libgomp-plugin-nvptx.so.1' dynamically links against 'libcuda.so.1', but testing is run on a system where there is no 'libcuda.so.1', this produces output such as: PASS: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 (test for excess errors) PASS: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 execution test FAIL: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 output pattern test, is CheCKpOInT libgomp: while loading libgomp-plugin-nvptx.so.1: libcuda.so.1: cannot open shared object file: No such file or directory ERROR STOP Error termination. Backtrace: [...] , should match CheCKpOInT( | |^M)+ERROR STOP ( | |^M)+Error termination.* ..., where after 'CheCKpOInT' we got 'libgomp: while loading [...]' injected before the expected 'ERROR STOP'. libgomp/ * testsuite/libgomp.oacc-fortran/error_stop-1.f: Initialize before the checkpoint. * 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.
2020-05-30openmp: omp_alloc(0, ...) should return NULL.Jakub Jelinek1-0/+25
2020-05-30 Jakub Jelinek <jakub@redhat.com> * allocator.c (omp_alloc): For size == 0, return NULL early. * testsuite/libgomp.c-c++-common/alloc-4.c: New test.
2020-05-23Fixes a hang on an invalid ID in a WAIT statement.Thomas Koenig1-0/+20
gcc/fortran/ChangeLog: 2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org> PR libfortran/95191 * libgfortran.h (libgfortran_error_codes): Add LIBERROR_BAD_WAIT_ID. libgfortran/ChangeLog: 2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org> PR libfortran/95191 * io/async.c (async_wait_id): Generate error if ID is higher than the highest current ID. * runtime/error.c (translate_error): Handle LIBERROR_BAD_WAIT_ID. libgomp/ChangeLog: 2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org> PR libfortran/95191 * testsuite/libgomp.fortran/async_io_9.f90: New test.
2020-05-19openmp: Add basic library allocator support.Jakub Jelinek3-0/+231
This patch adds very basic allocator support (omp_{init,destroy}_allocator, omp_{alloc,free}, omp_[sg]et_default_allocator). The plan is to use memkind (likely dlopened) for high bandwidth memory, but that part isn't implemented yet, probably mlock for pinned memory and see what other options there are for other kinds of memory. For offloading targets, we need to decide if we want to support the dynamic allocators (and on which targets), or if e.g. all we do is at compile time replace omp_alloc/omp_free calls with constexpr predefined allocators with something special. And allocate directive and allocator/uses_allocators clauses are future work too. 2020-05-19 Jakub Jelinek <jakub@redhat.com> * omp.h.in (omp_uintptr_t): New typedef. (__GOMP_UINTPTR_T_ENUM): Define. (omp_memspace_handle_t, omp_allocator_handle_t, omp_alloctrait_key_t, omp_alloctrait_value_t, omp_alloctrait_t): New typedefs. (__GOMP_DEFAULT_NULL_ALLOCATOR): Define. (omp_init_allocator, omp_destroy_allocator, omp_set_default_allocator, omp_get_default_allocator, omp_alloc, omp_free): Declare. * libgomp.h (struct gomp_team_state): Add def_allocator field. (gomp_def_allocator): Declare. * libgomp.map (OMP_5.0.1): Export omp_set_default_allocator, omp_get_default_allocator, omp_init_allocator, omp_destroy_allocator, omp_alloc and omp_free. * team.c (gomp_team_start): Copy over ts.def_allocator. * env.c (gomp_def_allocator): New variable. (parse_wait_policy): Adjust function comment. (parse_allocator): New function. (handle_omp_display_env): Print OMP_ALLOCATOR. (initialize_env): Call parse_allocator. * Makefile.am (libgomp_la_SOURCES): Add allocator.c. * allocator.c: New file. * icv.c (omp_set_default_allocator, omp_get_default_allocator): New functions. * testsuite/libgomp.c-c++-common/alloc-1.c: New test. * testsuite/libgomp.c-c++-common/alloc-2.c: New test. * testsuite/libgomp.c-c++-common/alloc-3.c: New test. * Makefile.in: Regenerated.
2020-05-14Add early return for invalid STATUS for close.Thomas Koenig1-0/+19
2020-05-14 Thomas Koenig <tkoenig@gcc.gnu.org> PR libfortran/95119 * io/close.c (close_status): Add CLOSE_INVALID. (st_close): Return early on invalid STATUS parameter. 2020-05-14 Thomas Koenig <tkoenig@gcc.gnu.org> PR libfortran/95119 * testsuite/libgomp.fortran/close_errors_1.f90: New test.
2020-05-14openmp: Also implicitly mark as declare target to functions mentioned in ↵Jakub Jelinek1-0/+51
target regions OpenMP 5.0 also specifies that functions referenced from target regions (except for target regions with device(ancestor:)) are also implicitly declare target to. This patch implements that. 2020-05-14 Jakub Jelinek <jakub@redhat.com> * function.h (struct function): Add has_omp_target bit. * omp-offload.c (omp_discover_declare_target_fn_r): New function, old renamed to ... (omp_discover_declare_target_tgt_fn_r): ... this. (omp_discover_declare_target_var_r): Call omp_discover_declare_target_tgt_fn_r instead of omp_discover_declare_target_fn_r. (omp_discover_implicit_declare_target): Also queue functions with has_omp_target bit set, for those walk with omp_discover_declare_target_fn_r, for declare target to functions walk with omp_discover_declare_target_tgt_fn_r. gcc/c/ * c-parser.c (c_parser_omp_target): Set cfun->has_omp_target. gcc/cp/ * cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target. gcc/fortran/ * trans-openmp.c: Include function.h. (gfc_trans_omp_target): Set cfun->has_omp_target. libgomp/ * testsuite/libgomp.c-c++-common/target-40.c: New test.
2020-05-13[Fortran] OpenMP - permit lastprivate in distribute + SIMD fixes (PR94690)Tobias Burnus7-0/+420
gcc/fortran/ 2020-05-13 Tobias Burnus <tobias@codesourcery.com> PR fortran/94690 * openmp.c (OMP_DISTRIBUTE_CLAUSES): Add OMP_CLAUSE_LASTPRIVATE. (gfc_resolve_do_iterator): Skip the private handling for SIMD as that is handled by ME code. * trans-openmp.c (gfc_trans_omp_do): Don't add private/lastprivate for dovar_found == 0, unless !simple. libgomp/ 2020-05-13 Tobias Burnus <tobias@codesourcery.com> PR fortran/94690 * testsuite/libgomp.fortran/pr66199-3.f90: New. * testsuite/libgomp.fortran/pr66199-4.f90: New. * testsuite/libgomp.fortran/pr66199-5.f90: New. * testsuite/libgomp.fortran/pr66199-6.f90: New. * testsuite/libgomp.fortran/pr66199-7.f90: New. * testsuite/libgomp.fortran/pr66199-8.f90: New. * testsuite/libgomp.fortran/pr66199-9.f90: New.
2020-05-12openmp: Implement discovery of implicit declare target to clausesJakub Jelinek1-0/+47
This attempts to implement what the OpenMP 5.0 spec in declare target section says as ammended by the 5.1 changes so far (related to device_type(host)), except that it doesn't have the device(ancestor: ...) handling yet because we do not support it yet, and I've left so far out the except lambda note, because I need that clarified. 2020-05-12 Jakub Jelinek <jakub@redhat.com> * omp-offload.h (omp_discover_implicit_declare_target): Declare. * omp-offload.c: Include context.h. (omp_declare_target_fn_p, omp_declare_target_var_p, omp_discover_declare_target_fn_r, omp_discover_declare_target_var_r, omp_discover_implicit_declare_target): New functions. * cgraphunit.c (analyze_functions): Call omp_discover_implicit_declare_target. * testsuite/libgomp.c/target-39.c: New test.
2020-04-29[gcn] Set 'UI_NONE' for 'TARGET_EXCEPT_UNWIND_INFO' [PR94282]Thomas Schwinge1-1/+0
In libgomp offloading testing, this resolves all the 'ld: error: undefined symbol: __gxx_personality_v0' FAILs. gcc/ PR target/94282 * common/config/gcn/gcn-common.c (gcn_except_unwind_info): New function. (TARGET_EXCEPT_UNWIND_INFO): Define. libgomp/ PR target/94282 * testsuite/libgomp.c-c++-common/function-not-offloaded.c: Remove 'dg-allow-blank-lines-in-output'.
2020-04-29[gcn] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE'Thomas Schwinge16-46/+46
..., per OpenACC 3.0, A.1.2. "AMD GPU Targets". This complements commit 6687d13a87c42dddc7d1c7adade38d31ba0d1401 "Rename acc_device_gcn to acc_device_radeon". libgomp/ * oacc-init.c (get_openacc_name): Handle 'gcn'. * testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type) [amdgcn*]: Return 'radeon'. Adjust all users. (check_effective_target_openacc_amdgcn_accel_present): Rename to... (check_effective_target_openacc_radeon_accel_present): ... this. Adjust all users. (check_effective_target_openacc_amdgcn_accel_selected): Rename to... (check_effective_target_openacc_radeon_accel_selected): ... this. Adjust all users.
2020-04-29Torture testing: 'libgomp.fortran/use_device_ptr-optional-2.f90'Thomas Schwinge1-0/+1
Fix-up for commit a2c26c50310a336361d8129ecdd43d3001d6cb3a (r278046) "Fortran] Support absent optional args with use_device_{ptr,addr}". libgomp/ * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Add 'dg-do run'.
2020-04-20Add 'dg-do run' to 'libgomp.fortran/target-enter-data-2.F90'Thomas Schwinge1-0/+1
Fix-up for commit af557050fd011a03d21dc26b31959033061a0443 "[OpenMP] Fix 'omp exit data' for Fortran arrays (PR 94635)". libgomp/ PR middle-end/94635 * testsuite/libgomp.fortran/target-enter-data-2.F90: Add 'dg-do run'.
2020-04-20Fix declare copyout in libgomp.oacc-c++/declare-pr94120.CTobias Burnus1-0/+1
Testing on the host does not make sense for 'declare copyout' for a same-scope stack-allocated variable. Once the copyout is done, the variable is gone. Hence, test the variable on the device. This can be revisit after the OpenACC semantic has been fixed; but with that fix, the test PASSes again with devices. PR middle-end/94120 * testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)' test case.
2020-04-17[OpenMP] Fix 'omp exit data' for Fortran arrays (PR 94635)Tobias Burnus1-0/+40
PR middle-end/94635 * gimplify.c (gimplify_scan_omp_clauses): Turn MAP_TO_PSET to MAP_DELETE. PR middle-end/94635 * testsuite/libgomp.fortran/target-enter-data-2.F90: New.
2020-04-13Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to ↵Thomas Schwinge20-23/+27
'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843] Fix-up for commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's "structured", not "static" data lifetimes/reference counters. libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c: ... this.
2020-04-10Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]Julian Brown16-0/+784
libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: New file. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c: Likewise.
generated by cgit v1.1 at 2025-09-25 18:50:48 +0000