aboutsummaryrefslogtreecommitdiff
path: root/libgomp
AgeCommit message (Collapse)AuthorFilesLines
2020-11-26Daily bump.GCC Administrator1-0/+10
2020-11-25Fix templatized C++ OpenACC 'cache' directive ICEsThomas Schwinge2-1/+24
This has been broken forever, whoops... gcc/cp/ * pt.c (tsubst_omp_clauses): Handle 'OMP_CLAUSE__CACHE_'. (tsubst_expr): Handle 'OACC_CACHE'. gcc/testsuite/ * c-c++-common/goacc/cache-1.c: Update. * c-c++-common/goacc/cache-2.c: Likewise. * g++.dg/goacc/cache-1.C: New. * g++.dg/goacc/cache-2.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c++/cache-1.C: New. * testsuite/libgomp.oacc-c-c++-common/cache-1.c: Update.
2020-11-25Fix atomic_capture-1.f90 testcaseAndrew Stubbs1-19/+19
The testcase had invalid assumptions about which loop iterations would run first and last. libgomp/ChangeLog * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 (main): Adjust expected results.
2020-11-25Daily bump.GCC Administrator1-0/+11
2020-11-24Tweak plugin-gcn.c definesAndrew Stubbs1-1/+2
Ensure the code will continue to compile when elf.h gets these definitions. libgomp/ChangeLog: * plugin/plugin-gcn.c: Don't redefine relocations if elf.h has them. (reserved): Delete unused define.
2020-11-24[testsuite] Avoid Tcl 8.5-specific behaviorThomas Schwinge2-0/+16
gcc/ * doc/install.texi (Prerequisites) <Tcl>: Add comment. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Reported-by: David Edelsohn <dje.gcc@gmail.com>
2020-11-19Daily bump.GCC Administrator1-0/+37
2020-11-18openmp: Retire nest-var ICV for OpenMP 5.1Kwok Cheung Yeung5-40/+90
This removes the nest-var ICV, expressing nesting in terms of the max-active-levels-var ICV instead. The max-active-levels-var ICV is now per data environment rather than per device. 2020-11-18 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_global_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_max_active_levels_var): Remove. (parse_boolean): Return true on success. (handle_omp_display_env): Express OMP_NESTED in terms of max_active_levels_var. Change format specifier for max_active_levels_var. (initialize_env): Set max_active_levels_var from OMP_MAX_ACTIVE_LEVELS, OMP_NESTED, OMP_NUM_THREADS and OMP_PROC_BIND. * icv.c (omp_set_nested): Express in terms of max_active_levels_var. (omp_get_nested): Likewise. (omp_set_max_active_levels): Use max_active_levels_var field instead of gomp_max_active_levels_var. (omp_get_max_active_levels): Likewise. * libgomp.h (struct gomp_task_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_supported_active_levels): Set to UCHAR_MAX. (gomp_max_active_levels_var): Delete. * libgomp.texi (omp_get_nested): Update documentation. (omp_set_nested): Likewise. (OMP_MAX_ACTIVE_LEVELS): Likewise. (OMP_NESTED): Likewise. (OMP_NUM_THREADS): Likewise. (OMP_PROC_BIND): Likewise. * parallel.c (gomp_resolve_num_threads): Replace reference to nest_var with max_active_levels_var. Use max_active_levels_var field instead of gomp_max_active_levels_var.
2020-11-18testsuite/libgomp.c/usleep.h: Use sleep-loop also for GCNTobias Burnus1-3/+4
As typically configured, newlib's libc.a does not build 'posix' and, hence, usleep is not available. Thus, use the same fallback as for nvptx. libgomp/ * testsuite/libgomp.c/usleep.h (fallback_usleep): Renamed from nvptx_usleep; use also for device={arch(gcn)}.
2020-11-15Daily bump.GCC Administrator1-0/+13
2020-11-14openmp: Add support for non-VLA {,first}private allocate on omp taskJakub Jelinek2-86/+195
This patch adds support for custom allocators on private/firstprivate clauses for task (and taskloop) constructs. Private didn't need anything special, but firstprivate if it is passed by reference needs the GOMP_alloc calls in the copyfn and GOMP_free in the task body. 2020-11-14 Jakub Jelinek <jakub@redhat.com> * gimplify.c (gimplify_omp_for): Add OMP_CLAUSE_ALLOCATE_ALLOCATOR decls as firstprivate on task clauses even when allocate clause decl is not lastprivate. * omp-low.c (install_var_field): Don't dereference omp_is_reference types if mask is 33 rather than 1. (scan_sharing_clauses): Populate allocate_map even for task constructs. For now remove it back for variables mentioned in reduction and in_reduction clauses on task/taskloop constructs or on VLA task firstprivates. For firstprivate on task construct, install the var field into field_map with by_ref and 33 instead of false and 1 if mentioned in allocate clause. (lower_private_allocate): Set TREE_THIS_NOTRAP on the created MEM_REF. (lower_rec_input_clauses): Handle allocate for task firstprivatized non-VLA variables. (create_task_copyfn): Likewise. * testsuite/libgomp.c-c++-common/allocate-1.c (struct S): New type. (foo): Add tests for non-VLA private and firstprivate clauses on omp task. (bar): Likewise. Remove taking of address from private/firstprivate variables. * testsuite/libgomp.c++/allocate-1.C (struct S): New type. (foo): Add p, q, px and s arguments. Add tests for array reductions and for non-VLA private and firstprivate clauses on omp task. (bar): Removed. (main): Adjust foo caller. Don't call bar.
2020-11-14Daily bump.GCC Administrator1-0/+24
2020-11-13Decompose OpenACC 'kernels' constructs into parts, a sequence of compute ↵Gergö Barany5-3/+66
constructs Not yet enabled by default: for now, the current mode of OpenACC 'kernels' constructs handling still remains '-fopenacc-kernels=parloops', but that is to change later. gcc/ * omp-oacc-kernels-decompose.cc: New. * Makefile.in (OBJS): Add it. * passes.def: Instantiate it. * tree-pass.h (make_pass_omp_oacc_kernels_decompose): Declare. * flag-types.h (enum openacc_kernels): Add. * doc/invoke.texi (-fopenacc-kernels): Document. * gimple.h (enum gf_mask): Add 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED', 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE', 'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'. (is_gimple_omp_oacc, is_gimple_omp_offloaded): Handle these. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * omp-expand.c (expand_omp_target, build_omp_regions_1) (omp_make_gimple_edges): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_for) (check_omp_nesting_restrictions, lower_oacc_reductions) (lower_oacc_head_mark, lower_omp_target): Likewise. * omp-offload.c (execute_oacc_device_lower): Likewise. gcc/c-family/ * c.opt (fopenacc-kernels): Add. gcc/fortran/ * lang.opt (fopenacc-kernels): Add. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: New. * c-c++-common/goacc/kernels-decompose-2.c: New. * c-c++-common/goacc/kernels-decompose-ice-1.c: New. * c-c++-common/goacc/kernels-decompose-ice-2.c: New. * gfortran.dg/goacc/kernels-decompose-1.f95: New. * gfortran.dg/goacc/kernels-decompose-2.f95: New. * c-c++-common/goacc/if-clause-2.c: Adjust. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: New. * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13Add 'libgomp.oacc-fortran/pr94358-1.f90' [PR94358]Gergö Barany1-0/+34
Document status quo re PR94358 "[OMP] Privatize internal array variables introduced by the Fortran FE". libgomp/ PR fortran/94358 * testsuite/libgomp.oacc-fortran/pr94358-1.f90: New. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13openmp: Support allocate for C/C++ array section reductionsJakub Jelinek1-6/+31
This adds allocate clause support for array section reductions. Furthermore, it fixes one bug that would cause inscan reductions with allocate to be rejected by C, and for now just ignores allocate for inscan/task reductions, that will need slightly more work. 2020-11-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.c (scan_sharing_clauses): For now remove for reduction clauses with inscan or task modifiers decl from allocate_map. (lower_private_allocate): Handle TYPE_P (new_var). (lower_rec_input_clauses): Handle allocate clause for C/C++ array reductions. gcc/c/ * c-typeck.c (c_finish_omp_clauses): Don't clear OMP_CLAUSE_REDUCTION_INSCAN unless reduction_seen == -2. libgomp/ * testsuite/libgomp.c-c++-common/allocate-1.c (foo): Add tests for array reductions. (main): Adjust foo callers.
2020-11-13Daily bump.GCC Administrator1-0/+18
2020-11-12openmp: Implement allocate clause in omp lowering.Jakub Jelinek6-6/+497
For now, task/taskloop constructs aren't handled and C/C++ array reductions and reductions with task or inscan modifiers need further work. Instead of calling omp_alloc/omp_free (where the former doesn't have alignment argument and omp_aligned_alloc is 5.1 only feature), this calls GOMP_alloc/GOMP_free, so that the library can fail if it would fall back into NULL (exception is zero length allocations). 2020-11-12 Jakub Jelinek <jakub@redhat.com> gcc/ * builtin-types.def (BT_FN_PTR_SIZE_SIZE_PTRMODE): New function type. * omp-builtins.def (BUILT_IN_GOACC_DECLARE): Move earlier. (BUILT_IN_GOMP_ALLOC, BUILT_IN_GOMP_FREE): New builtins. * gimplify.c (gimplify_scan_omp_clauses): Force allocator into a decl if it is not NULL, INTEGER_CST or decl. (gimplify_adjust_omp_clauses): Clear GOVD_EXPLICIT on explicit clauses which are being removed. Remove allocate clauses for variables not seen if they are private, firstprivate or linear too. Call omp_notice_variable on the allocator otherwise. (gimplify_omp_for): Handle iterator vars mentioned in allocate clauses similarly to non-is_gimple_reg iterators. * omp-low.c (struct omp_context): Add allocate_map field. (delete_omp_context): Delete it. (scan_sharing_clauses): Fill it from allocate clauses. Remove it if mentioned also in shared clause. (lower_private_allocate): New function. (lower_rec_input_clauses): Handle allocate clause for privatized variables, except for task/taskloop, C/C++ array reductions for now and task/inscan variables. (lower_send_shared_vars): Don't consider variables in allocate_map as shared. * omp-expand.c (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Use expand_omp_build_assign instead of gimple_build_assign + gsi_insert_after. * builtins.c (builtin_fnspec): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. * tree-ssa-ccp.c (evaluate_stmt): Handle BUILTIN_GOMP_ALLOC. * tree-ssa-dce.c (mark_stmt_if_obviously_necessary): Handle BUILTIN_GOMP_ALLOC. (mark_all_reaching_defs_necessary_1): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. (propagate_necessity): Likewise. gcc/fortran/ * f95-lang.c (ATTR_ALLOC_WARN_UNUSED_RESULT_SIZE_2_NOTHROW_LIST): Define. (gfc_init_builtin_functions): Add alloc_size and warn_unused_result attributes to __builtin_GOMP_alloc. * types.def (BT_PTRMODE): New primitive type. (BT_FN_VOID_PTR_PTRMODE, BT_FN_PTR_SIZE_SIZE_PTRMODE): New function types. libgomp/ * libgomp.map (GOMP_alloc, GOMP_free): Export at GOMP_5.0.1. * omp.h.in (omp_alloc): Add malloc and alloc_size attributes. * libgomp_g.h (GOMP_alloc, GOMP_free): Declare. * allocator.c (omp_aligned_alloc): New for now static function, add alignment argument and handle it. (omp_alloc): Reimplement using omp_aligned_alloc. (GOMP_alloc, GOMP_free): New functions. (omp_free): Add ialias. * testsuite/libgomp.c-c++-common/allocate-1.c: New test. * testsuite/libgomp.c++/allocate-1.C: New test.
2020-11-12Adjust 'libgomp.oacc-fortran/attach-descriptor-1.f90' for improved location ↵Thomas Schwinge1-2/+1
information Fix-up for commit b71ff8c15f5a7d6b1cc1524b4d27843f0d88dbda "Fortran: improve location data for OpenACC/OpenMP directives [PR97782]". libgomp/ PR fortran/97782 * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Adjust.
2020-11-11Daily bump.GCC Administrator1-0/+15
2020-11-10openmp: Implement OpenMP 5.0 base-pointer attachement and clause orderingChung-Lin Tang4-13/+122
This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang <cltang@codesourcery.com> gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
2020-11-06Daily bump.GCC Administrator1-0/+41
2020-11-05openmp: Mark deprecated symbols in OpenMP 5.0Kwok Cheung Yeung32-15/+77
2020-11-05 Ulrich Drepper <drepper@redhat.com> Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * Makefile.am (%.mod): Add -cpp and -fopenmp to compile flags. * Makefile.in: Regenerate. * fortran.c: Wrap uses of omp_set_nested and omp_get_nested with pragmas to ignore -Wdeprecated-declarations warnings. * icv.c: Likewise. * omp.h.in (__GOMP_DEPRECATED_5_0): Define. Mark omp_lock_hint_* enum values, omp_lock_hint_t, omp_set_nested, and omp_get_nested with __GOMP_DEPRECATED_5_0. * omp_lib.f90.in: Mark omp_get_nested and omp_set_nested as deprecated. * testsuite/libgomp.c++/affinity-1.C: Add -Wno-deprecated-declarations to test options. * testsuite/libgomp.c/affinity-1.c: Likewise. * testsuite/libgomp.c/affinity-2.c: Likewise. * testsuite/libgomp.c/appendix-a/a.15.1.c: Likewise. * testsuite/libgomp.c/lib-1.c: Likewise. * testsuite/libgomp.c/nested-1.c: Likewise. * testsuite/libgomp.c/nested-2.c: Likewise. * testsuite/libgomp.c/nested-3.c: Likewise. * testsuite/libgomp.c/pr32362-1.c: Likewise. * testsuite/libgomp.c/pr32362-2.c: Likewise. * testsuite/libgomp.c/pr32362-3.c: Likewise. * testsuite/libgomp.c/pr35549.c: Likewise. * testsuite/libgomp.c/pr42942.c: Likewise. * testsuite/libgomp.c/pr61200.c: Likewise. * testsuite/libgomp.c/sort-1.c: Likewise. * testsuite/libgomp.c/target-5.c: Likewise. * testsuite/libgomp.c/target-6.c: Likewise. * testsuite/libgomp.c/teams-1.c: Likewise. * testsuite/libgomp.c/thread-limit-1.c: Likewise. * testsuite/libgomp.c/thread-limit-2.c: Likewise. * testsuite/libgomp.c/thread-limit-4.c: Likewise. * testsuite/libgomp.fortran/affinity1.f90: Likewise. * testsuite/libgomp.fortran/lib1.f90: Likewise. * testsuite/libgomp.fortran/lib2.f: Likewise. * testsuite/libgomp.fortran/nested1.f90: Likewise. * testsuite/libgomp.fortran/teams1.f90: Likewise.
2020-11-02Daily bump.GCC Administrator1-0/+19
2020-11-02Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c' [PR85486]Thomas Schwinge3-97/+20
Avoid code duplication, and better test what we expect to happen. libgomp/ PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
2020-11-02libgomp testsuite: tell warning from error diagnostics, etc. [PR80219, PR85303]Thomas Schwinge1-0/+6
This changes makes 'dg-warning', 'dg-error', 'dg-bogus', 'dg-message' behave as expected, and also enables use of relative line numbers as well as 'dg-line'. libgomp/ PR testsuite/80219 PR testsuite/85303 * testsuite/lib/libgomp.exp (libgomp_init): Set 'gcc_warning_prefix', 'gcc_error_prefix'.
2020-10-30openmp: Use FIELD_TGT_EMPTY once moreJakub Jelinek1-1/+1
2020-10-30 Jakub Jelinek <jakub@redhat.com> * target.c (gomp_map_vars_internal): Use FIELD_TGT_EMPTY macro even in field_tgt_clear initializer.
2020-10-29Daily bump.GCC Administrator1-0/+19
2020-10-28openmp: Implicitly discover declare target for variants of declare variant callsJakub Jelinek1-0/+42
This marks all variants of declare variant also declare target if the base functions are called directly in target regions or declare target functions. 2020-10-28 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-offload.c (omp_declare_target_tgt_fn_r): Handle direct calls to declare variant base functions. libgomp/ * testsuite/libgomp.c/target-42.c: New test.
2020-10-28xfail and improve some failing libgomp tests [PR81690]Jakub Jelinek3-5/+31
With the patch I've posted today to fix up declare variant LTO handling, Tobias reported the patch still doesn't work, and there are two reasons for that. One is that when the base function is marked implicitly as declare target, we don't mark also implicitly the variants. I'll need to ask on omp-lang about details for that, but generally the compiler should do it some way. The other one is that the way base_delay is written, it will always call the usleep function, which is undesirable for nvptx. While the compiler will replace all direct calls to base_delay to nvptx_delay, the base_delay definition which calls usleep stays. 2020-10-28 Jakub Jelinek <jakub@redhat.com> Tom de Vries <tdevries@suse.de> PR testsuite/81690 * testsuite/libgomp.c/usleep.h: New file. * testsuite/libgomp.c/target-32.c: Include usleep.h. (main): Use tgt_usleep instead of usleep. * testsuite/libgomp.c/thread-limit-2.c: Include usleep.h. (main): Use tgt_usleep instead of usleep.
2020-10-28lto: LTO cgraph support for late declare variant resolution [PR96680]Jakub Jelinek1-0/+54
> I've tried to add the saving/restoring next to ipa refs saving/restoring, as > the declare variant alt stuff is kind of extension of those, unfortunately > following doesn't compile, because I need to also write or read a tree there > (ctx is a portion of DECL_ATTRIBUTES of the base function), but the ipa refs > write/read back functions don't have arguments that can be used for that. This patch adds the streaming out and in of those omp_declare_variant_alt hash table on the side data for the declare_variant_alt cgraph_nodes and treats for LTO purposes the declare_variant_alt nodes (which have no body) as if they contained a body that calls all the possible variants. After IPA all the calls to these magic declare_variant_alt calls are replaced with call to one of the variant depending on which one has the highest score in the context. 2020-10-28 Jakub Jelinek <jakub@redhat.com> PR lto/96680 gcc/ * lto-streamer.h (omp_lto_output_declare_variant_alt, omp_lto_input_declare_variant_alt): Declare variant. * symtab.c (symtab_node::get_partitioning_class): Return SYMBOL_DUPLICATE for declare_variant_alt nodes. * passes.c (ipa_write_summaries): Add declare_variant_alt to partition. * lto-cgraph.c (output_refs): Call omp_lto_output_declare_variant_alt on declare_variant_alt nodes. (input_refs): Call omp_lto_input_declare_variant_alt on declare_variant_alt nodes. * lto-streamer-out.c (output_function): Don't call collect_block_tree_leafs if DECL_INITIAL is error_mark_node. (lto_output): Call output_function even for declare_variant_alt nodes. * omp-general.c (omp_lto_output_declare_variant_alt, omp_lto_input_declare_variant_alt): New functions. gcc/lto/ * lto-common.c (lto_fixup_prevailing_decls): Don't use LTO_NO_PREVAIL on TREE_LIST's TREE_PURPOSE. * lto-partition.c (lto_balanced_map): Treat declare_variant_alt nodes like definitions. libgomp/ * testsuite/libgomp.c/declare-variant-1.c: New test.
2020-10-22Daily bump.GCC Administrator1-0/+31
2020-10-22openmp: Add test for OMP_TARGET_OFFLOAD=mandatory for cases where it must ↵Jakub Jelinek1-0/+33
not fail 2020-10-22 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/target-41.c: New test.
2020-10-22openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirementsJakub Jelinek7-21/+73
> Therefore, I think until omp_get_initial_device () value is changed, we The following so far untested patch implements that change. OpenMP 4.5 said for omp_get_initial_device: The value of the device number is implementation defined. If it is between 0 and one less than omp_get_num_devices() then it is valid for use with all device constructs and routines; if it is outside that range, then it is only valid for use with the device memory routines and not in the device clause. and OpenMP 5.0 similarly, but OpenMP 5.1 says: The value of the device number is the value returned by the omp_get_num_devices routine. As the new value is compatible with what has been required earlier, I think we can change it already now. 2020-10-22 Jakub Jelinek <jakub@redhat.com> * icv.c (omp_get_initial_device): Remove including corresponding ialias. * icv-device.c (omp_get_initial_device): New function. Return gomp_get_num_devices (). Add ialias. * target.c (resolve_device): Don't fail with OMP_TARGET_OFFLOAD=mandatory if device_id is equal to gomp_get_num_devices (). (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_pause_resource): Use gomp_get_num_devices () instead of GOMP_DEVICE_HOST_FALLBACK on the first use in the functions, in uses dominated by the gomp_get_num_devices call use num_devices_openmp instead. * libgomp.texi (omp_get_initial_device): Document. * config/gcn/icv-device.c (omp_get_initial_device): New function. Add ialias. * config/nvptx/icv-device.c (omp_get_initial_device): Likewise. * testsuite/libgomp.c/target-40.c: New test.
2020-10-21libgomp: Hopefully avoid false positive warnings in env.c on solarisJakub Jelinek1-6/+2
> the patch also breaks bootstrap on both i386-pc-solaris2.11 and > sparc-sun-solaris2.11: > > /vol/gcc/src/hg/master/local/libgomp/env.c: In function 'initialize_env': > /vol/gcc/src/hg/master/local/libgomp/env.c:414:16: error: 'new_offload' may be used uninitialized in this function [-Werror=maybe-uninitialized] > 414 | *offload = new_offload; > | ~~~~~~~~~^~~~~~~~~~~~~ > /vol/gcc/src/hg/master/local/libgomp/env.c:384:30: note: 'new_offload' was declared here > 384 | enum gomp_target_offload_t new_offload; > | ^~~~~~~~~~~ I can't reproduce that, but I fail to see why we need two separate variables, one with actual value and one tracking if the value is valid. So, I'm going with: 2020-10-21 Jakub Jelinek <jakub@redhat.com> * env.c (parse_target_offload): Change new_offload var type to int, preinitialize to -1, remove found var and test new_offload != -1 instead of found.
2020-10-21Daily bump.GCC Administrator1-0/+27
2020-10-20libgomp: Fix up bootstrap in libgomp/target.c due to false positive warningJakub Jelinek1-37/+39
> On 10/20/20 2:11 PM, Tobias Burnus wrote: > > > Unfortunately, the committed patch > > (r11-4121-g1bfc07d150790fae93184a79a7cce897655cb37b) > > causes build errors. > > > > The error seems to be provoked by function cloning – as the code > > itself looks fine: > > ... > > struct gomp_device_descr *devices_s > > = malloc (num_devices * sizeof (struct gomp_device_descr)); > > ... > > for (i = 0; i < num_devices; i++) > > if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) > > devices_s[num_devices_after_openmp++] = devices[i]; > > gomp_target_init.part.0 () > { > ... > <bb 2> > devices_s_1 = malloc (0); > ... > num_devices.16_67 = num_devices; > ... > if (num_devices.16_67 > 0) > goto <bb 3>; [89.00%] > else > goto <bb 18>; [11.00%] > > Which seems to have an ordering problem. This patch fixes the warning that breaks the bootstrap. 2020-10-20 Jakub Jelinek <jakub@redhat.com> * target.c (gomp_target_init): Inside of the function, use automatic variables corresponding to num_devices, num_devices_openmp and devices global variables and update the globals only at the end of the function.
2020-10-20openmp: Implement support for OMP_TARGET_OFFLOAD environment variableKwok Cheung Yeung4-8/+130
This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled. It may be set to MANDATORY (abort if offloading cannot be performed), DISABLED (no offloading to devices) or DEFAULT (offload to device if possible, fall back to host if not). 2020-10-20 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_target_offload_var): New. (parse_target_offload): New. (handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD. (initialize_env): Parse OMP_TARGET_OFFLOAD. * libgomp.h (gomp_target_offload_t): New. (gomp_target_offload_var): New. * libgomp.texi (OMP_TARGET_OFFLOAD): New section. * target.c (resolve_device): Generate error if device not found and offloading is mandatory. (gomp_target_fallback): Generate error if offloading is mandatory. (GOMP_target): Add argument in call to gomp_target_fallback. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Generate error if offloading is mandatory. (GOMP_target_data): Add argument in call to gomp_target_data_fallback. (GOMP_target_data_ext): Likewise. (gomp_target_task_fn): Add argument in call to gomp_target_fallback. (gomp_target_init): Return early if offloading is disabled.
2020-10-16Daily bump.GCC Administrator1-0/+5
2020-10-15libgomp: Amend documentation for omp_get_max_active_levels and ↵Kwok Cheung Yeung1-3/+3
omp_get_supported_active_levels 2020-10-15 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * libgomp.texi (omp_get_max_active_levels): Modify description. (omp_get_supported_active_levels): Make descriptions consistent.
2020-10-15Daily bump.GCC Administrator1-0/+4
2020-10-14libgomp: Fix a typo in documentationJakub Jelinek1-1/+1
2020-10-14 Jakub Jelinek <jakub@redhat.com> * libgomp.texi (omp_get_supported_active_levels): Fix a typo.
2020-10-14Daily bump.GCC Administrator1-0/+24
2020-10-13openmp: Add support for the omp_get_supported_active_levels runtime library ↵Kwok Cheung Yeung11-4/+70
routine This patch implements the omp_get_supported_active_levels runtime routine from the OpenMP 5.0 specification, which returns the maximum number of active nested parallel regions supported by this implementation. The current maximum (set using the omp_set_max_active_levels routine or the OMP_MAX_ACTIVE_LEVELS environment variable) cannot exceed this number. 2020-10-13 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_max_active_levels_var): Initialize to gomp_supported_active_levels. (initialize_env): Limit gomp_max_active_levels_var to be at most equal to gomp_supported_active_levels. * fortran.c (omp_get_supported_active_levels): Add ialias_redirect. (omp_get_supported_active_levels_): New. * icv.c (omp_set_max_active_levels): Limit gomp_max_active_levels_var to at most equal to gomp_supported_active_levels. (omp_get_supported_active_levels): New. * libgomp.h (gomp_supported_active_levels): New. * libgomp.map (OMP_5.0.1): Add omp_get_supported_active_levels and omp_get_supported_active_levels_. * libgomp.texi (omp_get_supported_active_levels): New. (omp_set_max_active_levels): Update. Add reference to omp_get_supported_active_levels. * omp.h.in (omp_get_supported_active_levels): New. * omp_lib.f90.in (omp_get_supported_active_levels): New. * omp_lib.h.in (omp_get_supported_active_levels): New. * testsuite/libgomp.c/lib-2.c (main): Check omp_get_max_active_levels against omp_get_supported_active_levels. * testsuite/libgomp.fortran/lib4.f90 (lib4): Likewise.
2020-10-12Daily bump.GCC Administrator1-0/+5
2020-10-11aix: remove libgomp and libatomic archives before creating FAT archivesClément Chigot1-0/+3
AIX caches shared objects in archives with read-other permission. libgomp and libatomic might be in use during the build or testing, which may cause archiver operations on them to fail. This patch adjusts the Makefile fragments to delete the library archives before creating fresh archives containing both the 32 bit and 64 bit shared objects. libatomic/ChangeLog: 2020-10-11 Clement Chigot <clement.chigot@atos.net> * config/t-aix: Delete and recreate libatomic before creating FAT library. libgomp/ChangeLog: 2020-10-11 Clement Chigot <clement.chigot@atos.net> * config/t-aix: Delete and recreate libgomp before creating FAT library.
2020-10-09Daily bump.GCC Administrator1-0/+6
2020-10-08[libgomp, nvptx] Report launch dimensions in GOMP_OFFLOAD_runTom de Vries1-1/+8
Using this patch, when using GOMP_DEBUG=1 and launching a kernel in GOMP_OFFLOAD_run (used by the omp implementation), we see the kernel launch dimensions: ... GOMP_OFFLOAD_run: kernel main$_omp_fn$0: \ launch [(teams: 1), 1, 1] [(lanes: 32), (threads: 1), 1] ... Build on x86_64-linux with nvptx accelerator, tested libgomp. libgomp/ChangeLog: 2020-10-08 Tom de Vries <tdevries@suse.de> PR libgomp/81802 * plugin/plugin-nvptx.c (GOMP_OFFLOAD_run): Report launch dimensions.
2020-10-07Daily bump.GCC Administrator1-0/+9
2020-10-06[openacc, libgomp, testsuite] Xfail declare-5.f90Tom de Vries1-0/+1
We're currently running into: ... FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -fomit-frame-pointer \ -funroll-loops -fpeel-loops -ftracer -finline-functions execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -g execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -Os execution test ... A PR was filed for this: PR92790 - "[OpenACC] declare device_resident - Fortran common blocks not handled / libgomp.oacc-fortran/declare-5.f90 fails" Xfail the fails. Tested on x86_64-linux with nvptx accelerator. libgomp/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-fortran/declare-5.f90: Add xfail for PR92790.
2020-10-06[openacc] Fix acc declare for VLAsTom de Vries1-5/+0
Consider test-case test.c, with VLA A: ... int main (void) { int N = 1000; int A[N]; #pragma acc declare copy(A) return 0; } ... compiled using: ... $ gcc test.c -fopenacc -S -fdump-tree-all ... At original, we have: ... #pragma acc declare map(tofrom:A); ... but at gimple, we have a map (to:A.1), but not a map (from:A.1): ... int[0:D.2074] * A.1; { int A[0:D.2074] [value-expr: *A.1]; saved_stack.2 = __builtin_stack_save (); try { A.1 = __builtin_alloca_with_align (D.2078, 32); #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076]) } finally { __builtin_stack_restore (saved_stack.2); } } ... This is caused by the following incompatibility. When storing the desired from clause in oacc_declare_returns, we use 'A.1' as the key: ... 10898 oacc_declare_returns->put (decl, c); (gdb) call debug_generic_expr (decl) A.1 (gdb) call debug_generic_expr (c) map(from:(*A.1)) ... but when looking it up, we use 'A' as the key: ... (gdb) 1471 tree *c = oacc_declare_returns->get (t); (gdb) call debug_generic_expr (t) A ... Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr. In addition, unshare the looked up value, to fix avoid running into an "incorrect sharing of tree nodes" error. Using these two fixes, we get our desired: ... finally { + #pragma omp target oacc_declare map(from:(*A.1)) __builtin_stack_restore (saved_stack.2); } ... Build on x86_64-linux with nvptx accelerator, tested libgomp. gcc/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> PR middle-end/90861 * gimplify.c (gimplify_bind_expr): Handle lookup in oacc_declare_returns using key with decl-expr. libgomp/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> PR middle-end/90861 * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.
generated by cgit v1.1 at 2024-09-27 13:57:52 +0000