]> git.ipfire.org Git - thirdparty/gcc.git/log
thirdparty/gcc.git
5 years agoFix testsuite failures due to expected compiler warnings
Kwok Cheung Yeung [Mon, 2 Mar 2020 16:32:08 +0000 (08:32 -0800)] 
Fix testsuite failures due to expected compiler warnings

2020-03-02  Kwok Cheung Yeung  <kcy@codesourcery.com>

libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (main): Add
dg-warnings for expected compiler warnings.
* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c (DoWorkVec):
Likewise.
* testsuite/libgomp.oacc-fortran/optional-private.f90 (t1, t2):
Likewise.

5 years agoFix ICE in oacc_verify_routine_clauses
Kwok Cheung Yeung [Thu, 20 Feb 2020 20:56:13 +0000 (12:56 -0800)] 
Fix ICE in oacc_verify_routine_clauses

2020-02-20  Kwok Cheung Yeung  <kcy@codesourcery.com>

gcc/
* omp-general.c (oacc_verify_routine_clauses): Handle
OMP_CLAUSE_NOHOST.

5 years agolibgomp - fix dg-warning line numbers
Tobias Burnus [Wed, 18 Sep 2019 11:45:34 +0000 (13:45 +0200)] 
libgomp - fix dg-warning line numbers

libgomp/
        * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Fix dg-warning
        line numbers.

5 years agoFix OpenACC "ephemeral" asynchronous host-to-device copies
Julian Brown [Wed, 11 Sep 2019 20:22:03 +0000 (13:22 -0700)] 
Fix OpenACC "ephemeral" asynchronous host-to-device copies

libgomp/
* libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_host2dev): Update
prototype.
* libgomp.h (gomp_copy_host2dev): Update prototype.
* oacc-host.c (host_openacc_async_host2dev): Add ephemeral parameter.
* oacc-mem.c (memcpy_tofrom_device): Update call to gomp_copy_host2dev.
(update_dev_host): Likewise.
* oacc-parallel.c (GOACC_enter_exit_data): Call async versions of
acc_attach/acc_detach/acc_detach_finalize functions.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_async_host2dev): Add
ephemeral parameter.  Copy source data to temporary space immediately
if true, and pass to queue_push_copy.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_openacc_async_host2dev):
Add EPHEMERAL parameter, and FIXME function comment.
* target.c (goacc_device_copy_async): Remove.
(gomp_copy_host2dev): Add ephemeral parameter. Update function comment.
Call async host2dev plugin hook directly.
(gomp_copy_dev2host): Call async dev2host plugin hook directly.
(gomp_map_vars_existing, gomp_map_pointer, gomp_attach_pointer,
gomp_detach_pointer): Update calls to gomp_copy_host2dev.
(gomp_map_vars_internal): Don't use coalescing buffer for asynchronous
copies. Update calls to gomp_copy_host2dev.
(gomp_update): Update calls to gomp_copy_host2dev.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c (main): Fix
async-safety issue. Increase number of iterations.
* testsuite/libgomp.oacc-fortran/lib-16.f90: Fix async-safety issue.
* testsuite/libgomp.oacc-fortran/lib-16-2.f90: Likewise.

5 years agoFix warning syntax and typos in two libgomp tests
Julian Brown [Mon, 20 May 2019 23:31:41 +0000 (16:31 -0700)] 
Fix warning syntax and typos in two libgomp tests

libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Expect
"optimized:" not "note:" in warnings.

5 years agoFortran "declare create"/allocate support for OpenACC
Julian Brown [Tue, 26 Feb 2019 23:48:00 +0000 (15:48 -0800)] 
Fortran "declare create"/allocate support for OpenACC

2018-10-04  Cesar Philippidis  <cesar@codesourcery.com>
            Julian Brown  <julian@codesourcery.com>

gcc/
* omp-low.c (scan_sharing_clauses): Update handling of OpenACC declare
create, declare copyin and declare deviceptr to have local lifetimes.
(convert_to_firstprivate_int): Handle pointer types.
(convert_from_firstprivate_int): Likewise.  Create local storage for
the values being pointed to.  Add new orig_type argument.
(lower_omp_target): Handle GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
Add orig_type argument to convert_from_firstprivate_int call.
Allow pointer types with GOMP_MAP_FIRSTPRIVATE_INT.  Don't privatize
firstprivate VLAs.
* tree-pretty-print.c (dump_omp_clause): Handle
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.

gcc/fortran/
* gfortran.h (enum gfc_omp_map_op): Add OMP_MAP_DECLARE_ALLOCATE,
OMP_MAP_DECLARE_DEALLOCATE.
(gfc_omp_clauses): Add update_allocatable.
* trans-array.c (gfc_array_allocate): Call
gfc_trans_oacc_declare_allocate for decls that have oacc_declare_create
attribute set.
* trans-decl.c (find_module_oacc_declare_clauses): Relax
oacc_declare_create to OMP_MAP_ALLOC, and oacc_declare_copyin to
OMP_MAP_TO, in order to match OpenACC 2.5 semantics.
* trans-openmp.c (gfc_trans_omp_clauses): Use GOMP_MAP_ALWAYS_POINTER
(for update directive) or GOMP_MAP_FIRSTPRIVATE_POINTER (otherwise) for
allocatable scalar decls.  Handle OMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}
clauses.
(gfc_trans_oacc_executable_directive): Use GOMP_MAP_ALWAYS_POINTER
for allocatable scalar data clauses inside acc update directives.
(gfc_trans_oacc_declare_allocate): New function.
* trans-stmt.c (gfc_trans_allocate): Call
gfc_trans_oacc_declare_allocate for decls with oacc_declare_create
attribute set.
(gfc_trans_deallocate): Likewise.
* trans.h (gfc_trans_oacc_declare_allocate): Declare.

gcc/testsuite/
* gfortran.dg/goacc/declare-allocatable-1.f90: New test.

include/
* gomp-constants.h (enum gomp_map_kind): Define
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE} and GOMP_MAP_FLAG_SPECIAL_4.

libgomp/
* oacc-mem.c (gomp_acc_declare_allocate): New function.
* oacc-parallel.c (GOACC_enter_exit_data): Handle
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
* testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-2.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-4.f90: New test.

2020-02-19  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
* trans-openmp.c (gfc_omp_check_optional_argument): Handle non-decl
case.

gcc/
* gimplify.c (gimplify_scan_omp_clauses): Handle
GOMP_MAP_DECLARE_ALLOCATE and GOMP_MAP_DECLARE_DEALLOCATE.

libgomp/
* libgomp.h (gomp_acc_declare_allocate): Remove prototype.
* oacc-mem.c (gomp_acc_declare_allocate): Make static.  Add POINTER
argument. Use acc_delete instead of acc_free.  Handle scalar
mappings.
(find_group_last): Handle GOMP_MAP_DECLARE_ALLOCATE and
GOMP_MAP_DECLARE_DEALLOCATE groupings.
(goacc_enter_data_internal): Fix kind check for
GOMP_MAP_DECLARE_ALLOCATE. Pass new pointer argument to
gomp_acc_declare_allocate.
(goacc_exit_data_internal): Unlock device mutex around
gomp_acc_declare_allocate call. Pass new pointer argument. Handle
group pointer mapping for deallocate.

5 years agoRe-do OpenACC private variable resolution
Julian Brown [Mon, 14 Oct 2019 20:12:39 +0000 (13:12 -0700)] 
Re-do OpenACC private variable resolution

gcc/
* config/gcn/gcn-protos.h (gcn_goacc_adjust_gangprivate_decl): Rename
to...
(gcn_goacc_adjust_private_decl): ...this.
* config/gcn/gcn-tree.c (diagnostic-core.h): Include.
(gcn_goacc_adjust_gangprivate_decl): Rename to...
(gcn_goacc_adjust_private_decl): ...this. Add LEVEL parameter.
* config/gcn/gcn.c (TARGET_GOACC_ADJUST_GANGPRIVATE_DECL): Rename to...
(TARGET_GOACC_ADJUST_PRIVATE_DECL): ...this.
* config/nvptx/nvptx.c (tree-pretty-print.h): Include.
(nvptx_goacc_adjust_private_decl): New function.
(TARGET_GOACC_ADJUST_PRIVATE_DECL): Define hook using above function.
* doc/tm.texi.in (TARGET_GOACC_ADJUST_GANGPRIVATE_DECL): Rename to...
(TARGET_GOACC_ADJUST_PRIVATE_DECL): ...this.
* doc/tm.texi: Regenerated.
* internal-fn.c (expand_UNIQUE): Handle IFN_UNIQUE_OACC_PRIVATE.
* internal-fn.h (IFN_UNIQUE_CODES): Add OACC_PRIVATE.
* omp-low.c (omp_context): Remove oacc_partitioning_levels field.
(lower_oacc_reductions): Add PRIVATE_MARKER parameter.  Insert before
fork.
(lower_oacc_head_tail): Add PRIVATE_MARKER parameter. Modify its
gimple call arguments as appropriate. Don't set
oacc_partitioning_levels in omp_context. Pass private_marker to
lower_oacc_reductions.
(oacc_record_private_var_clauses): Don't check for NULL ctx.
(mark_oacc_gangprivate): Remove unused function.
(make_oacc_private_marker): New function.
(lower_omp_for): Only call oacc_record_vars_in_bind for
OpenACC contexts.  Create private marker and pass to
lower_oacc_head_tail.
(lower_omp_target): Remove unnecessary call to
oacc_record_private_var_clauses. Remove call to mark_oacc_gangprivate.
Create private marker and pass to lower_oacc_reductions.
(process_oacc_gangprivate_1): Remove.
(lower_omp_1): Only call oacc_record_vars_in_bind for OpenACC.  Don't
iterate over contexts calling process_oacc_gangprivate_1.
(omp-offload.c (oacc_loop_xform_head_tail): Treat
private-variable markers like fork/join when transforming head/tail
sequences.
(execute_oacc_device_lower): Use IFN_UNIQUE_OACC_PRIVATE instead of
"oacc gangprivate" attributes to determine partitioning level of
variables.  Remove unused variables.
* omp-sese.c (find_gangprivate_vars): New function.
(find_local_vars_to_propagate): Use GANGPRIVATE_VARS parameter instead
of "oacc gangprivate" attribute to determine which variables are
gang-private.
(oacc_do_neutering): Use find_gangprivate_vars.
* target.def (adjust_gangprivate_decl): Rename to...
(adjust_private_decl): ...this.  Update documentation (briefly).

libgomp/
* testsuite/libgomp.oacc-fortran/gangprivate-attrib-1.f90: Use
oaccdevlow dump and update scanned output.
* testsuite/libgomp.oacc-fortran/gangprivate-attrib-2.f90: Likewise.
Add missing atomic to force worker partitioning for test variable.

5 years agoHandle references in OpenACC "private" clauses
Julian Brown [Fri, 20 Sep 2019 20:53:10 +0000 (13:53 -0700)] 
Handle references in OpenACC "private" clauses

gcc/
* gimplify.c (localize_reductions): Rewrite references for
OMP_CLAUSE_PRIVATE also.

libgomp/
* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: New test.
* testsuite/libgomp.oacc-c++/privatized-ref-2.C: New test.
* testsuite/libgomp.oacc-c++/privatized-ref-3.C: New test.

5 years agoReduce testsuite fails
Tobias Burnus [Thu, 19 Sep 2019 13:57:08 +0000 (15:57 +0200)] 
Reduce testsuite fails

gcc/testsuite/
2019-09-19  Tobias Burnus  <tobias@codesourcery.com>

        * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add
        one dg-message for additional -fopt-info-optimized-omp output.
        * gfortran.dg/goacc/classify-kernels.f95: Likewise.
        * gfortran.dg/goacc/kernels-decompose-1.f95: Change 'note' to
        'optimized' in dg-message.

5 years agoSilence compiler warnings
Tobias Burnus [Wed, 18 Sep 2019 06:44:20 +0000 (08:44 +0200)] 
Silence compiler warnings

gcc/
2019-09-17  Tobias Burnus  <tobias@codesourcery.com>

        * config/gcn/gcn.c (gcn_expand_scalar_to_vector_address,
        gcn_md_reorg): Remove unused statement.
        (gcn_emutls_var_init): Add missing return - after sorry abort.
        * config/gcn/gcn.md (movdi_symbol_save_scc): Fix condition.
        * config/gcn/mkoffload.c (process_obj): Remove unused variables.
        * gimplify.c (gomp_oacc_needs_data_present): Likewise.
        (gimplify_adjust_omp_clauses): Fix condition by adding ().
        * omp-low.c (process_oacc_gangprivate_1): Comment unused
        parameter name to silence unused warning.
        * omp-sese.c (omp_sese_number, omp_sese_pseudo): Remove
        superfluous ().
        (oacc_do_neutering): Use signed int to avoid a warning.
        * tree-ssa-structalias.c (find_func_aliases_for_builtin_call,
        find_func_clobbers): Use unsigned to silence warning.

5 years agoOpenACC profiling-interface fixes for asynchronous operations
Julian Brown [Wed, 11 Sep 2019 03:34:45 +0000 (20:34 -0700)] 
OpenACC profiling-interface fixes for asynchronous operations

libgomp/
* oacc-host.c (host_openacc_async_queue_callback): Invoke callback
function immediately.
* oacc-parallel.c (struct async_prof_callback_info, async_prof_dispatch,
queue_async_prof_dispatch): New.
(GOACC_parallel_keyed): Call queue_async_prof_dispatch for asynchronous
profile-event dispatches.
(GOACC_update): Likewise.
* oacc-mem.c (GOACC_enter_exit_data): Call queue_async_prof_dispatch
for asynchronous profile-event dispatches.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
(cb_compute_construct_start): Remove/fix TODO.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
(cb_exit_data_start): Tweak expected state values.
(cb_exit_data_end): Likewise.
(cb_compute_construct_start): Remove/fix TODO.
(cb_compute_construct_end): Don't do adjustments for
acc_ev_enqueue_launch_start/acc_ev_enqueue_launch_end callbacks.
(cb_compute_construct_end): Tweak expected state values.
(cb_enqueue_launch_start, cb_enqueue_launch_end): Don't expect
launch-enqueue operations to happen synchronously with respect to
profiling events on async streams.
(main): Tweak expected state values.
* testsuite/libgomp.oacc-c-c++-common/lib-94.c (main): Reorder
operations for async-safety.

5 years agoUpdate expected messages, errors and warnings for "kernels" tests
Julian Brown [Fri, 13 Sep 2019 01:03:17 +0000 (18:03 -0700)] 
Update expected messages, errors and warnings for "kernels" tests

gcc/testsuite/
* c-c++-common/goacc/classify-kernels-unparallelized.c: Update expected
message/warning/error output.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/kernels-decompose.c: Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c: Likewise.
* c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loop-auto.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-4-extern.c: Likewise.

5 years agoA couple of GCN-specific test fixes
Julian Brown [Wed, 11 Sep 2019 15:31:38 +0000 (08:31 -0700)] 
A couple of GCN-specific test fixes

libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (check): Skip
vector dimension test for AMD GCN.

5 years agoUse more appropriate var in localize_reductions call
Julian Brown [Fri, 6 Sep 2019 15:46:20 +0000 (08:46 -0700)] 
Use more appropriate var in localize_reductions call

gcc/
* gimplify.c (gimplify_omp_for): Use for_stmt in call to
localize_reductions.

5 years agoRemove duplicate SESE code in NVPTX backend
Julian Brown [Fri, 6 Sep 2019 11:53:17 +0000 (04:53 -0700)] 
Remove duplicate SESE code in NVPTX backend

gcc/
* config/nvptx/nvptx.c (omp-sese.h): Include.
(bb_pair_t, bb_pair_vec_t, pseudo_node_t, bracket, bracket_vec_t,
bb_sese, bb_sese::~bb_sese, bb_sese::append, bb_sese::remove,
BB_SET_SESE, BB_GET_SESE, nvptx_sese_number, nvptx_sese_pseudo,
nvptx_sese_color, nvptx_find_sese): Remove.
(nvptx_neuter_pars): Call omp_find_sese instead of nvptx_find_sese.
* omp-sese.c (omp-sese.h): Include.
(struct parallel): Rename to...
(struct parallel_g): This.
(parallel::parallel, parallel::~parallel): Rename to...
(parallel_g::parallel_g, parallel_g::~parallel_g): These.
(omp_sese_dump_pars, omp_sese_find_par, omp_sese_discover_pars,
populate_single_mode_bitmaps, find_ssa_names_to_propagate,
find_partitioned_var_uses, find_local_vars_to_propagate,
neuter_worker_single): Update for parallel_g name change.
(bb_pair_t, bb_pair_vec_t): Remove.
(omp_find_sese): Make global.
* omp-sese.h (bb_pair_t, bb_pair_vec_t): New.
(omp_find_sese): Add prototype.

5 years agoFix tree check failure with reduction localization
Julian Brown [Fri, 6 Sep 2019 11:42:16 +0000 (04:42 -0700)] 
Fix tree check failure with reduction localization

gcc/
* gimplify.c (gimplify_omp_workshare): Use OMP_CLAUSES, OMP_BODY
instead of OMP_TARGET_CLAUSES, OMP_TARGET_BODY.

5 years agoEnable worker partitioning for AMD GCN
Julian Brown [Tue, 3 Sep 2019 15:57:39 +0000 (08:57 -0700)] 
Enable worker partitioning for AMD GCN

gcc/
* config/gcn/gcn.c (gcn_goacc_validate_dims): Remove
no-flag_worker-partitioning assertion.
(TARGET_GOACC_WORKER_PARTITIONING): Define target hook to true.
* config/gcn/gcn.opt (flag_worker_partitioning): Change default to 1.

libgomp/
* plugin/plugin-gcn.c (gcn_exec): Change default number of workers to
16.

5 years agoReference reduction localization
Julian Brown [Tue, 3 Sep 2019 15:54:28 +0000 (08:54 -0700)] 
Reference reduction localization

gcc/
* gimplify.c (privatize_reduction): New struct.
(localize_reductions_r, localize_reductions): New functions.
(gimplify_omp_for): Call localize_reductions.
(gimplify_omp_workshare): Likewise.
* omp-low.c (lower_oacc_reductions): Handle localized reductions.
Create fewer temp vars.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE_REDUCTION_PRIVATE_DECL
documentation.
* tree.c (omp_clause_num_ops): Bump number of ops for
OMP_CLAUSE_REDUCTION to 6.
(walk_tree_1): Adjust accordingly.
* tree.h (OMP_CLAUSE_REDUCTION_PRIVATE_DECL): Add macro.

5 years agoFix up tests for oaccdevlow pass splitting
Julian Brown [Tue, 3 Sep 2019 20:37:50 +0000 (13:37 -0700)] 
Fix up tests for oaccdevlow pass splitting

gcc/testsuite/
* c-c++-common/goacc/classify-kernels-unparallelized.c,
c-c++-common/goacc/classify-kernels.c,
c-c++-common/goacc/classify-parallel.c,
c-c++-common/goacc/classify-routine.c,
gfortran.dg/goacc/classify-kernels-unparallelized.f95,
gfortran.dg/goacc/classify-kernels.f95,
gfortran.dg/goacc/classify-parallel.f95,
gfortran.dg/goacc/classify-routine.f95: Scan oaccloops dump instead of
oaccdevlow pass.

5 years agoAMD GCN adjustments for middle-end worker partitioning
Julian Brown [Wed, 4 Sep 2019 23:33:02 +0000 (16:33 -0700)] 
AMD GCN adjustments for middle-end worker partitioning

gcc/
* config/gcn/gcn-protos.h (gcn_goacc_adjust_propagation_record): Rename
prototype to...
(gcn_goacc_create_propagation_record): This.
* config/gcn/gcn-tree.c (gcn_goacc_adjust_propagation_record): Rename
function to...
(gcn_goacc_create_propagation_record): This.  Adjust comment.
* config/gcn/gcn.c (gcn_init_builtins): Override decls for
        BUILT_IN_GOACC_SINGLE_START, BUILT_IN_GOACC_SINGLE_COPY_START,
        BUILT_IN_GOACC_SINGLE_COPY_END and BUILT_IN_GOACC_BARRIER.
(gcn_fork_join): Remove inaccurate comment.
(TARGET_GOACC_ADJUST_PROPAGATION_RECORD): Rename to...
(TARGET_GOACC_CREATE_PROPAGATION_RECORD): This.

5 years agoOpenACC middle-end worker-partitioning support
Julian Brown [Wed, 11 Oct 2017 15:07:18 +0000 (08:07 -0700)] 
OpenACC middle-end worker-partitioning support

gcc/
* Makefile.in (OBJS): Add omp-sese.o.
* omp-builtins.def (BUILT_IN_GOACC_BARRIER, BUILT_IN_GOACC_SINGLE_START,
        BUILT_IN_GOACC_SINGLE_COPY_START, BUILT_IN_GOACC_SINGLE_COPY_END): New
builtins.
* omp-offload.c (omp-sese.h): Include header.
(oacc_loop_xform_head_tail): Call update_stmt for modified builtin
calls.
(oacc_loop_process): Likewise.
(default_goacc_create_propagation_record): New default implementation
for TARGET_GOACC_CREATE_PROPAGATION_RECORD hook.
(execute_oacc_loop_designation): New.  Split out of oacc_device_lower.
(execute_oacc_gimple_workers): New.  Likewise.
(execute_oacc_device_lower): Recreate dims array.
(pass_data_oacc_loop_designation, pass_data_oacc_gimple_workers): New.
(pass_oacc_loop_designation, pass_oacc_gimple_workers): New.
(make_pass_oacc_loop_designation, make_pass_oacc_gimple_workers): New.
* omp-offload.h (oacc_fn_attrib_level): Add prototype.
* omp-sese.c: New file.
* omp-sese.h: New file.
* passes.def (pass_oacc_loop_designation, pass_oacc_gimple_workers):
Add passes.
* target.def (worker_partitioning, create_propagation_record): Add
target hooks.
* targhooks.h (default_goacc_create_propagation_record): Add prototype.
* tree-pass.h (make_pass_oacc_loop_designation,
make_pass_oacc_gimple_workers): Add prototypes.
* doc/tm.texi.in (TARGET_GOACC_WORKER_PARTITIONING,
TARGET_GOACC_CREATE_PROPAGATION_RECORD): Add documentation hooks.
* doc/tm.texi: Regenerate.

5 years agoTarget-dependent gang-private variable decl rewriting
Julian Brown [Tue, 3 Sep 2019 23:35:10 +0000 (16:35 -0700)] 
Target-dependent gang-private variable decl rewriting

gcc/
* omp-offload.c (convert.h): Include.
(struct addr_expr_rewrite_info): Add struct.
(rewrite_addr_expr): New function.
(is_sync_builtin_call): New function.
(execute_oacc_device_lower): Support rewriting gang-private variables
using target hook, and fix up addr_expr nodes afterwards.
* target.def (adjust_gangprivate_decl): New target hook.
* doc/tm.texi.in (TARGET_GOACC_ADJUST_GANGPRIVATE_DECL): Document new
target hook.
* doc/tm.texi: Regenerate.

5 years agoFix libgomp.oacc-fortran/lib-13.f90 async bug
Julian Brown [Tue, 3 Sep 2019 14:57:05 +0000 (07:57 -0700)] 
Fix libgomp.oacc-fortran/lib-13.f90 async bug

libgomp/
* testsuite/libgomp.oacc-fortran/lib-13.f90: End data region after
wait API calls.

5 years agoWait at end of OpenACC asynchronous kernels regions
Julian Brown [Fri, 9 Aug 2019 20:01:33 +0000 (13:01 -0700)] 
Wait at end of OpenACC asynchronous kernels regions

gcc/
* omp-oacc-kernels.c (add_wait): New function, split out of...
(add_async_clauses_and_wait): ...here. Call new outlined function.
(decompose_kernels_region_body): Add wait at the end of
explicitly-asynchronous kernels regions.

5 years agoAdd missing exec_params libgomp plugin entry points
Julian Brown [Mon, 5 Aug 2019 22:05:35 +0000 (15:05 -0700)] 
Add missing exec_params libgomp plugin entry points

libgomp/
* plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_exec_params,
GOMP_OFFLOAD_openacc_async_exec_params): New functions.

5 years agoNVPTX GOMP_OFFLOAD_openacc_async_construct arg fix and gomp_print_* support
Julian Brown [Mon, 29 Jul 2019 22:05:35 +0000 (15:05 -0700)] 
NVPTX GOMP_OFFLOAD_openacc_async_construct arg fix and gomp_print_* support

libgomp/
* config/nvptx/gomp_print.c (gomp_print_string, gomp_print_integer,
gomp_print_double): New.

5 years agoMake OpenACC function-parameter explosion optional
Julian Brown [Fri, 26 Jul 2019 20:51:48 +0000 (13:51 -0700)] 
Make OpenACC function-parameter explosion optional

* configure.ac (amdgcn*-*-*): Add target-libffi to noconfigdirs for AMD
GCN.
* configure: Regenerated.

gcc/
* builtin-types.def (BT_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
Remove.
* config/i386/i386.c (ix86_goacc_explode_args): New.
(TARGET_GOACC_EXPLODE_ARGS): Define, using above function.
* doc/tm.texi: Regenerated.
* doc/tm.texi.in: Add TARGET_GOACC_EXPLODE_ARGS hook.
* fortran/types.def (BT_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
Remove.
* omp-builtins.def (GOACC_parallel_keyed_v2): Remove.
* omp-expand.c (expand_omp_target): Use explode_args target hook.
Use GOMP_LAUNCH_ARGS_EXPLODED launch tag.  Handle exploded args for
serial target kinds.
* omp-low.c (build_receiver_ref, build_sender_ref,
create_omp_child_function, scan_omp_target, lower_omp_target): Use
explode_args target hook.
* target.def (explode_args): New target hook.
* tree-ssa-structalias.c (target.h): Include.
(find_func_aliases_for_builtin_call): Conditionalise disabling of pass
for OpenACC parallel regions based on explode_args target hook.  Remove
'params' from BUILT_IN_GOACC_PARALLEL arguments.
(find_func_clobbers): Likewise.
(ipa_pta_execute): Update for removed 'params' argument.

include/
* gomp-constants.h (GOMP_LAUNCH_ARGS_EXPLODED): Define.

libgomp/
* libgomp.map (GOMP_2.0.GOMP_4_BRANCH): Remove GOACC_parallel_keyed_v2.
* libgomp_g.h (GOACC_parallel_keyed_v2): Remove prototype.
* oacc-parallel.c (GOACC_parallel_keyed_internal): Rename to...
(GOACC_parallel_keyed): ...this.  Handle GOMP_LAUNCH_ARGS_EXPLODED
launch tag.  Remove previous wrapper functions.
(GOACC_parallel_keyed_v2): Remove.

5 years agoAdd changes to profiling interface from OG8 branch
Kwok Cheung Yeung [Fri, 21 Jun 2019 17:40:38 +0000 (10:40 -0700)] 
Add changes to profiling interface from OG8 branch

This bundles up the parts of the profiling code from the OG8 branch that were
not included in the upstream patch.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Update.

libgomp/
* oacc-init.c (get_property_any): Add profiling code.

libgomp/
* Makefile.am (libgomp_la_SOURCES): Add
oacc-profiling-acc_register_library.c.
* Makefile.in: Regenerate.
* libgomp.texi: Remove paragraph about acc_register_library.
* oacc-parallel.c (GOACC_parallel_keyed_internal): Set device_api for
profiling.
* oacc-profiling-acc_register_library.c: New file.
* oacc-profiling.c (goacc_profiling_initialize): Call
acc_register_library.  Avoid duplicate registration.
(acc_register_library): Remove.
* config/nvptx/oacc-profiling-acc_register_library.c:
New empty file.
* config/nvptx/oacc-profiling.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-dispatch-1.c: Remove
call to acc_register_library.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-valid_bytes-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-version-1.c: Likewise.

5 years agoAllow the accelerator to have more offloaded functions than the host
Cesar Philippidis [Sun, 7 Jul 2019 18:25:51 +0000 (11:25 -0700)] 
Allow the accelerator to have more offloaded functions than the host

libgomp/
* target.c (gomp_load_image_to_device): Allow the accelerator to
possess more offloaded functions than the host.

5 years agoAssumed-size arrays with non-lexical data mappings
Julian Brown [Fri, 5 Jul 2019 01:14:41 +0000 (18:14 -0700)] 
Assumed-size arrays with non-lexical data mappings

gcc/
* gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for
assumed-size arrays in map clauses for Fortran/OpenMP.
* omp-low.c (lower_omp_target): Set the size of assumed-size Fortran
arrays to one to allow use of data already mapped on the offload device.

gcc/fortran/
* trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping
assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type.

5 years agoApply gangprivate attribute to innermost decl
Julian Brown [Tue, 28 May 2019 15:42:10 +0000 (08:42 -0700)] 
Apply gangprivate attribute to innermost decl

...and fix parallelism-level calculation when applying the attribute.

gcc/
* omp-low.c (mark_oacc_gangprivate): Add CTX parameter.  Use to look up
correct decl to add attribute to.
(lower_omp_for): Move "oacc gangprivate" processing from here...
(process_oacc_gangprivate_1): ...to here. New function.
(lower_omp_target): Update call to mark_oacc_gangprivate.
(execute_lower_omp): Call process_oacc_gangprivate_1 for each OMP
context.

libgomp/
* testsuite/libgomp.oacc-fortran/gangprivate-attrib-2.f90: New test.

5 years agoFix expected messages in goacc tests
Kwok Cheung Yeung [Fri, 31 May 2019 19:25:03 +0000 (12:25 -0700)] 
Fix expected messages in goacc tests

The expected messages in the OpenACC kernel-related tests should be prefixed
with 'optimized:' rather than 'note:'.

2019-05-31  Kwok Cheung Yeung  <kcy@codesourcery.com>

gcc/testsuite/
* c-c++-common/goacc/kernels-decompose-1.c: Change 'note:' to
'optimized:'.  Fix typo.
* c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c:
Change 'note:' to 'optimized:'.
* c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c: Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c:
Likewise.
* c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c:
Likewise.
* c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loop-auto.c: Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise.

5 years agoFix lexically-nested data mappings for no_alloc or optional arguments
Julian Brown [Tue, 21 May 2019 00:27:38 +0000 (17:27 -0700)] 
Fix lexically-nested data mappings for no_alloc or optional arguments

gcc/
* gimplify.c (gimplify_adjust_omp_clauses_1): Support implied no_alloc
and optional arguments based on mappings in enclosing data regions.

5 years agoFix references declared in lexically-enclosing OpenACC data region
Julian Brown [Sun, 19 May 2019 17:42:20 +0000 (10:42 -0700)] 
Fix references declared in lexically-enclosing OpenACC data region

gcc/fortran/
* trans-openmp.c (gfc_omp_finish_clause): Guard addition of clauses for
pointers with DECL_P.

gcc/
* gimplify.c (oacc_array_mapping_info): Add REF field.
(gimplify_scan_omp_clauses): Initialise above field for data blocks
passed by reference.
(gomp_oacc_needs_data_present): Handle references.
(gimplify_adjust_omp_clauses_1): Handle references and optional
arguments for variables declared in lexically-enclosing OpenACC data
region.

5 years agoAdd kernels for-index reuse testcase.
Julian Brown [Thu, 16 May 2019 12:47:16 +0000 (05:47 -0700)] 
Add kernels for-index reuse testcase.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-for-index-reuse-1.c: New
test.

5 years agoAvoid introducing 'create' mapping clauses for loop index variables in kernels regions
Julian Brown [Thu, 16 May 2019 12:45:35 +0000 (05:45 -0700)] 
Avoid introducing 'create' mapping clauses for loop index variables in kernels regions

gcc/
* omp-oacc-kernels.c (find_omp_for_index_vars_1,
find_omp_for_index_vars): New functions.
(maybe_build_inner_data_region): Add IDX_VARS argument. Don't add
CREATE mapping clauses for loop index variables.  Set TREE_ADDRESSABLE
flag on newly-mapped declarations as a side effect.
(decompose_kernels_region_body): Call find_omp_for_index_vars.  Don't
create PRESENT clause for loop index variables.  Pass index variable
set to maybe_build_inner_data_region.

5 years agoUpdate OpenACC version to 2.6
Julian Brown [Wed, 9 Jan 2019 11:41:04 +0000 (03:41 -0800)] 
Update OpenACC version to 2.6

gcc/c-family/
* c-cppbuiltin.c (c_cpp_builtins): Update _OPENACC define to 201711.

gcc/doc/
* invoke.texi: Update mention of OpenACC version to 2.6.

gcc/fortran/
* cpp.c (cpp_define_builtins): Update _OPENACC define to 201711.
* gfortran.texi: Update mentions of OpenACC version to 2.6.
* intrinsic.texi: Likewise.

gcc/testsuite/
* c-c++-common/cpp/openacc-define-3.c: Update expected value for
_OPENACC define.
* gfortran.dg/openacc-define-3.f90: Likewise.

libgomp/
* libgomp.texi: Update mentions of OpenACC version to 2.6.  Update
section numbers to match version 2.6 of the spec.
* openacc.f90 (openacc_version): Update to 201711.
* openacc_lib.h (openacc_version): Update to 201711.
* testsuite/libgomp.oacc-fortran/openacc_version-1.f: Update expected
openacc_version to 201711.
* testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.

5 years agoAdjust parallelism of loops in gang-single parts of OpenACC kernels regions: "struct...
Thomas Schwinge [Fri, 1 Feb 2019 17:12:05 +0000 (18:12 +0100)] 
Adjust parallelism of loops in gang-single parts of OpenACC kernels regions: "struct adjust_nested_loop_clauses_wi_info"

The current code apparently is too freaky at least for for GCC 4.6:

    [...]/gcc/omp-oacc-kernels.c: In function 'tree_node* transform_kernels_loop_clauses(gimple*, tree, tree, tree, tree)':
    [...]/gcc/omp-oacc-kernels.c:584:10: error: expected identifier before numeric constant
    [...]/gcc/omp-oacc-kernels.c: In lambda function:
    [...]/gcc/omp-oacc-kernels.c:584:25: error: expected '{' before '=' token
    [...]/gcc/omp-oacc-kernels.c: In function 'tree_node* transform_kernels_loop_clauses(gimple*, tree, tree, tree, tree)':
    [...]/gcc/omp-oacc-kernels.c:584:25: warning: lambda expressions only available with -std=c++0x or -std=gnu++0x [enabled by default]
    [...]/gcc/omp-oacc-kernels.c:584:28: error: no match for 'operator=' in '{} = & loop_gang_clause'
    [...]

gcc/
* omp-oacc-kernels.c (struct adjust_nested_loop_clauses_wi_info): New.
(adjust_nested_loop_clauses, transform_kernels_loop_clauses): Use it.

5 years agoMake new OpenACC kernels conversion the default; adjust and add tests
Thomas Schwinge [Wed, 23 Jan 2019 10:40:08 +0000 (02:40 -0800)] 
Make new OpenACC kernels conversion the default; adjust and add tests

gcc/c-family/
* c.opt (fopenacc-kernels): Default to "split".
gcc/fortran/
* lang.opt (fopenacc-kernels): Default to "split".
gcc/
* doc/invoke.texi (-fopenacc-kernels): Update.
gcc/testsuite/
* c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c:
New file.
* c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c:
Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise.
* c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c:
Likewise.
* c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c:
Likewise.
* c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loop-auto.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c:
Likewise.
* c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise.
* c-c++-common/goacc/classify-kernels-unparallelized.c: Update.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* c-c++-common/goacc/dtype-1.c: Likewise.
* c-c++-common/goacc/if-clause-2.c: Likewise.
* c-c++-common/goacc/kernels-conversion.c: Likewise.
* c-c++-common/goacc/kernels-decompose-1.c: Likewise.
* c-c++-common/goacc/loop-2-kernels.c: Likewise.
* c-c++-common/goacc/note-parallelism.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/dtype-1.f95: Likewise.
* gfortran.dg/goacc/kernels-conversion.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
Update.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90:
Likewise.

5 years agoNew OpenACC kernels region decompose algorithm
Thomas Schwinge [Thu, 24 Jan 2019 16:40:03 +0000 (08:40 -0800)] 
New OpenACC kernels region decompose algorithm

Previously, OpenACC kernels region bodies were decomposed into a sequence of
alternating gang-single and gang-parallel "parallel" regions. The new
algorithm in this patch introduces a third possibility: Loops that look like
they might benefit from the parloops pass are converted into old "kernels"
regions, exposing them to the parloops pass later on. This has the benefit
that loops that cannot be parallelized are not offloaded to the GPU.

gcc/
* omp-oacc-kernels.c (adjust_region_code_walk_stmt_fn)
(adjust_region_code): New functions.
(make_loops_gang_single): Update.
(make_gang_single_region): Rename to...
(make_region_seq): ... this, and update.
(make_gang_parallel_loop_region): Rename to...
(make_region_loop_nest): ... this, and update.
(is_unconditional_oacc_for_loop): Remove stmt parameter and check.
(decompose_kernels_region_body): Update.
gcc/testsuite/
* c-c++-common/goacc/kernels-conversion.c: Adjust test.
* gfortran.dg/goacc/kernels-conversion.f95: Likewise.
* c-c++-common/goacc/kernels-decompose-1.c: New file.
* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: New
file.

5 years agoLaunch kernels asynchronously in OpenACC kernels regions
Gergö Barany [Mon, 21 Jan 2019 20:50:14 +0000 (12:50 -0800)] 
Launch kernels asynchronously in OpenACC kernels regions

Kernels regions are decomposed into one or more smaller regions that are to
be executed in sequence. With this patch, all of these regions are launched
asynchronously, and a wait directive is added after them. This means that
the host only waits once for the kernels to complete, not once per kernel.
If the original kernels region was marked async, that asynchronous behavior
is preserved, and no wait is added.

    gcc/
    * omp-oacc-kernels.c (add_async_clauses_and_wait): New function...
    (decompose_kernels_region_body): ... called from here.

    gcc/testsuite/
    * c-c++-common/goacc/kernels-conversion.c: Test automatically generated
    async clauses.
    * gfortran.dg/goacc/kernels-conversion.f95: Likewise.

5 years agoAdjust parallelism of loops in gang-single parts of OpenACC kernels regions
Gergö Barany [Thu, 24 Jan 2019 06:11:11 +0000 (22:11 -0800)] 
Adjust parallelism of loops in gang-single parts of OpenACC kernels regions

Loops in gang-single parts of kernels regions cannot be executed in
gang-redundant mode. If the user specified gang clauses on such loops, emit
an error and remove these clauses. Adjust automatic partitioning to exclude
gang partitioning in gang-single regions.

    gcc/
    * omp-oacc-kernels.c (add_parent_or_loop_num_clause): New function.
    (adjust_nested_loop_clauses): Likewise.
    (transform_kernels_loop_clauses, make_gang_parallel_loop_region):
    Add worker and vector clause parameters, emit error on illegal
    nesting.
    (visit_loops_in_gang_single_region): Emit warning on conditionally
    executed code with a gang clause.
    (make_loops_gang_single): New function.
    (decompose_kernels_region_body): Separate out gang/worker/vector clauses
    for separate handling; add call to make_loops_gang_single.
    * omp-offload.c (oacc_loop_auto_partitions): Add and propagate
    is_oacc_gang_single parameter.
    (oacc_loop_partition): Likewise.
    (execute_oacc_device_lower): Adjust call to oacc_loop_partition.

5 years agoHandle conditional execution of loops in OpenACC kernels regions
Gergö Barany [Wed, 23 Jan 2019 22:32:57 +0000 (14:32 -0800)] 
Handle conditional execution of loops in OpenACC kernels regions

Any OpenACC loop controlled by an if statement or a non-OpenACC loop must be
executed in a gang-single region. Detecting such loops is not trivial as
OpenACC kernels expansion is done on GIMPLE but before computation of the
control flow graph. This patch adds an auxiliary analysis for determining
whether a statement is inside a conditionally executed region (relative to
the kernels region's entry).

    gcc/
    * omp-oacc-kernels.c (control_flow_regions): New class.
    (control_flow_regions::control_flow_regions): New constructor.
    (control_flow_regions::is_unconditional_oacc_for_loop): New method.
    (control_flow_regions::find_rep): Likewise.
    (control_flow_regions::union_reps): Likewise.
    (control_flow_regions::compute_regions): Likewise.
    (decompose_kernels_region_body): Use test for conditional execution.

    gcc/testsuite/
    * c-c++-common/goacc/kernels-conversion.c: Add test for conditionally
    executed code.
    * gfortran.dg/goacc/kernels-conversion.f95: Likewise.

5 years agoTurn OpenACC kernels regions into a sequence of parallel regions
Gergö Barany [Mon, 21 Jan 2019 15:16:06 +0000 (07:16 -0800)] 
Turn OpenACC kernels regions into a sequence of parallel regions

This patch decomposes each OpenACC kernels region into a sequence of
parallel regions. Each OpenACC loop nest turns into its own region; any code
between such loop nests is gathered up into a region as well. The loop
regions can be distributed across gangs if the original kernels region had a
num_gangs clause, while the other regions are executed in "gang-single"
mode. The implied default "auto" clause on kernels loops is made explicit
unless there is a conflicting clause.

    gcc/
    * omp-oacc-kernels.c (top_level_omp_for_in_stmt): New function.
    (make_gang_single_region): Likewise.
    (transform_kernels_loop_clauses, make_gang_parallel_loop_region):
    Likewise.
    (flatten_binds): Likewise.
    (make_data_region_try_statement): Likewise.
    (maybe_build_inner_data_region): Likewise.
    (decompose_kernels_region_body): Likewise.
    (transform_kernels_region): Delegate to decompose_kernels_region_body
    and make_data_region_try_statement.

    gcc/testsuite/
    * c-c++-common/goacc/kernels-conversion.c: Test for a gang-single
    region.
    * gfortran.dg/goacc/kernels-conversion.f95: Likewise.

5 years agoSeparate OpenACC kernels regions in data and parallel parts
Gergö Barany [Mon, 21 Jan 2019 13:28:20 +0000 (05:28 -0800)] 
Separate OpenACC kernels regions in data and parallel parts

This is the first in a series of patches that completely rework the handling
of the OpenACC "kernels" directive. In the future, kernels regions will be
transformed into data regions containing a sequence of serial and parallel
offloaded regions. This first patch sets up a new pass that is responsible
for this transformation, and in a first step constructs the new data region
containing a parallel region with the original kernels region's body.

gcc/
* Makefile.in: Add...
* omp-oacc-kernels.c: ... this new file for the kernels conversion
pass.
* flag-types.h (enum openacc_kernels): Add "split" style.  Adjust
all users.
* doc/invoke.texi (-fopenacc-kernels): Update.
* passes.def: Add pass_convert_oacc_kernels to pipeline.
* tree-pass.h (make_pass_convert_oacc_kernels): Add declaration.
gcc/testsuite/
* c-c++-common/goacc/kernels-conversion.c: New test.
* gfortran.dg/goacc/kernels-conversion.f95: Likewise.
* c-c++-common/goacc/if-clause-2.c: Update.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.

5 years agoAdd OpenACC target kinds for decomposed kernels regions
Thomas Schwinge [Wed, 23 Jan 2019 14:56:52 +0000 (06:56 -0800)] 
Add OpenACC target kinds for decomposed kernels regions

This patch is in preparation for changes that will cut up OpenACC kernels
regions into individual parts. For the new sub-regions that will be
generated, this adds the following new kinds of OpenACC regions for internal
use:

- GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED for parts of kernels
  regions to be executed in gang-redundant mode
- GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE for parts of kernels
  regions to be executed in gang-single mode
- GF_OMP_TARGET_KIND_OACC_DATA_KERNELS for data regions generated around the
  body of a kernels region

    gcc/
    * gimple.h (enum gf_mask): Add new target kinds
    GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED,
    GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE, and
    GF_OMP_TARGET_KIND_OACC_DATA_KERNELS.
    (is_gimple_omp_oacc): Handle new target kinds.
    (is_gimple_omp_offloaded): Likewise.
    * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
    * omp-expand.c (expand_omp_target): Likewise.
    (build_omp_regions_1): Likewise.
    (omp_make_gimple_edges): Likewise.
    * omp-low.c (is_oacc_parallel_or_serial): Likewise.
    (was_originally_oacc_kernels): New function.
    (scan_omp_for): Update check for illegal nesting.
    (check_omp_nesting_restrictions): Handle new target kinds.
    (lower_oacc_reductions): Likewise.
    (lower_omp_target): Likewise.
    * omp-offload.c (execute_oacc_device_lower): Likewise.

5 years agoUse "-fopenacc-kernels=parloops" to document "parloops" test cases
Thomas Schwinge [Wed, 30 Jan 2019 09:32:10 +0000 (10:32 +0100)] 
Use "-fopenacc-kernels=parloops" to document "parloops" test cases

gcc/
* flag-types.h (enum openacc_kernels): New type.
gcc/c-family/
* c.opt (fopenacc-kernels): New flag.
gcc/fortran/
* lang.opt (fopenacc-kernels): New flag.
gcc/testsuite/
* c-c++-common/goacc/kernels-1.c: Add
"-fopenacc-kernels=parloops".
* c-c++-common/goacc/kernels-acc-loop-reduction.c: Likewise.
* c-c++-common/goacc/kernels-acc-loop-smaller-equal.c: Likewise.
* c-c++-common/goacc/kernels-alias-2.c: Likewise.
* c-c++-common/goacc/kernels-alias-3.c: Likewise.
* c-c++-common/goacc/kernels-alias-4.c: Likewise.
* c-c++-common/goacc/kernels-alias-5.c: Likewise.
* c-c++-common/goacc/kernels-alias-6.c: Likewise.
* c-c++-common/goacc/kernels-alias-7.c: Likewise.
* c-c++-common/goacc/kernels-alias-8.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta-2.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta-4.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta.c: Likewise.
* c-c++-common/goacc/kernels-alias.c: Likewise.
* c-c++-common/goacc/kernels-counter-var-redundant-load.c:
Likewise.
* c-c++-common/goacc/kernels-counter-vars-function-scope.c:
Likewise.
* c-c++-common/goacc/kernels-double-reduction-n.c: Likewise.
* c-c++-common/goacc/kernels-double-reduction.c: Likewise.
* c-c++-common/goacc/kernels-loop-2-acc-loop.c: Likewise.
* c-c++-common/goacc/kernels-loop-2.c: Likewise.
* c-c++-common/goacc/kernels-loop-3-acc-loop.c: Likewise.
* c-c++-common/goacc/kernels-loop-3.c: Likewise.
* c-c++-common/goacc/kernels-loop-acc-loop.c: Likewise.
* c-c++-common/goacc/kernels-loop-data-2.c: Likewise.
* c-c++-common/goacc/kernels-loop-data-enter-exit-2.c: Likewise.
* c-c++-common/goacc/kernels-loop-data-enter-exit.c: Likewise.
* c-c++-common/goacc/kernels-loop-data-update.c: Likewise.
* c-c++-common/goacc/kernels-loop-data.c: Likewise.
* c-c++-common/goacc/kernels-loop-g.c: Likewise.
* c-c++-common/goacc/kernels-loop-mod-not-zero.c: Likewise.
* c-c++-common/goacc/kernels-loop-n-acc-loop.c: Likewise.
* c-c++-common/goacc/kernels-loop-n.c: Likewise.
* c-c++-common/goacc/kernels-loop-nest.c: Likewise.
* c-c++-common/goacc/kernels-loop.c: Likewise.
* c-c++-common/goacc/kernels-one-counter-var.c: Likewise.
* c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c:
Likewise.
* c-c++-common/goacc/kernels-reduction.c: Likewise.
* gfortran.dg/goacc/kernels-alias-2.f95: Likewise.
* gfortran.dg/goacc/kernels-alias-3.f95: Likewise.
* gfortran.dg/goacc/kernels-alias-4.f95: Likewise.
* gfortran.dg/goacc/kernels-alias.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-2.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-data.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-inner.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-n.f95: Likewise.
* gfortran.dg/goacc/kernels-loop.f95: Likewise.
* gfortran.dg/goacc/kernels-loops-adjacent.f95: Likewise.
* gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95:
Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c:
Add "-fopenacc-kernels=parloops".
* testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-update.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90:
Likewise.

5 years agoUse functional parameters for data mappings in OpenACC child functions
Cesar Philippidis [Thu, 21 Dec 2017 21:40:34 +0000 (13:40 -0800)] 
Use functional parameters for data mappings in OpenACC child functions

* Makefile.def: Make libgomp depend on libffi.
* configure.ac: Likewise.
* Makefile.in: Regenerate.
* configure: Regenerate.

gcc/fortran/
* types.def: (BF_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
Define.

gcc/
* builtin-types.def (BF_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
Define.
* config/nvptx/nvptx.c (nvptx_expand_cmp_swap): Handle PARM_DECLs.
* omp-builtins.def (BUILD_IN_GOACC_PARALLEL): Call
GOACC_parallel_keyed_v2.
* omp-expand.c (expand_omp_target): Update call to
BUILT_IN_GOACC_PARALLEL.
* omp-low.c (struct omp_context): Add parm_map member.
(lookup_parm): New function.
(build_receiver_ref): Lookup parm_map decls.
(install_parm_decl): New function.
(install_var_field): Install parm_map decl for OpenACC parallel region
data clauses.
(delete_omp_context): Clean parm_map.
(scan_sharing_clauses): Install subarray variable mapping into parm_map.
(create_omp_child_function): Defer creation of child function for
OpenACC parallel regions.
(scan_omp_target): Likewise.
(append_decl_arg): New function.
(lower_omp_target): Create an child offloaded function using one
parameter per data mapping for OpenACC parallel regions.
* tree-ssa-structalias.c (find_func_aliases_for_builtin_call):
Ignore OpenACC parallel regions.
(find_func_clobbers): Likewise.
(ipa_pta_execute): Likewise.

libgomp/
* Makefile.am: Add libffi build dependency.
* configure.ac: Likewise.
* Makefile.in: Regenerate.
* config.h.in: Regenerate.
* configure: Regenerate.
* libgomp-plugin.h: Define GOMP_OFFLOAD_openacc_exec_params and
GOMP_OFFLOAD_openacc_async_exec_params.
* libgomp.h (acc_dispatch_t): Use them here.
* libgomp.map (GOACC_parallel_keyed_v2): Declare.
* libgomp_g.h (GOACC_parallel_keyed_v2): Likewise.
* oacc-host.c (host_openacc_exec_params): New function.
(host_openacc_async_exec_params): Likewise.
* oacc-parallel.c (goacc_call_host_fn): Likewise.
(GOACC_parallel_keyed_internal): Likewise.
(GOACC_parallel_keyed): Wrapper for GOACC_parallel_keyed_internal.
(GOACC_parallel_keyed_v2): Likewise.
* plugin/plugin-nvptx.c (nvptx_exec): Replace CUDeviceptr dp parameter
with void **kargs.
(openacc_exec_internal): New function.
(GOMP_OFFLOAD_openacc_exec_params): New function.
(GOMP_OFFLOAD_openacc_exec): Update to call openacc_exec_internal.
(openacc_async_exec_internal): New function.
(GOMP_OFFLOAD_openacc_async_exec_params): New function.
(GOMP_OFFLOAD_openacc_async_exec): Update call to
openacc_async_exec_internal.
* target.c (gomp_load_plugin_for_device): Handle
openacc_exec_params and openacc_async_exec_params.
* testsuite/Makefile.in: Regenerate.
* testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
Xfail on offloaded targets.

* Makefile.def: Bootstrap module libffi. Add libffi dependency
to all-target-libgomp.
* Makefile.in: Regenerate.
* configure.ac: Add libffi to bootstrap_target_libs when libgomp
is bootstrapped.
* configure: Regenerate.

gcc/
* omp-low.c (install_parm_decl): Don't extract identifiers from
artifical decls.

gcc/testsuite/
* c-c++-common/goacc/large_array.c: New test.

(cherry picked from openacc-gcc-7-branch commit
b4dd21b9a1f9f499c613b55225cad689b7928a7f, commit
9ba1d875dcb9412cccdd49138a3525e7adab3e76, commit
762cf3c7890fab15a69494a6480455cd99621d7d, and commit
6585af7290fd79f6cb834a39c2bbf7e1934808b1)

5 years agoDisable libstdc++ dependency for libffi
Maciej W. Rozycki [Thu, 20 Dec 2018 14:10:16 +0000 (14:10 +0000)] 
Disable libstdc++ dependency for libffi

Disable AC_PROG_CXX and consequently a libstdc++ dependency for libffi,
introduced with upstream libffi commit 7d698125b1f0 ("Use the proper C++
compiler to run C++ tests").  This is only needed for the libffi test
suite, which we don't have to support in the GCC tree, as libffi is
maintained as a separate project.  The dependency causes a build failure
with the `powerpc64le-linux-gnu' target due to a circular dependency:

make[1]: Circular configure-target-libffi <- maybe-all-target-libstdc++-v3 dependency dropped.
make[1]: *** [configure-target-libffi] Error 1
make: *** [all] Error 2

due to a libgomp dependency for libstdc++ and then a libffi dependency
for libgomp, introduced with commit 998eb38b265d ("Use functional
parameters for data mappings in OpenACC child functions").

/
* Makefile.def (lang_env_dependencies): Disable `cxx' dependency
for `libffi'.
* Makefile.in: Regenerate.

libffi/
* configure.ac: Disable AC_PROG_CXX.
* configure: Regenerate.
* Makefile.in: Regenerate.
* include/Makefile.in: Regenerate.
* man/Makefile.in: Regenerate.
* testsuite/Makefile.in: Regenerate.

5 years agoAdd support for gang local storage allocation in shared memory
Julian Brown [Thu, 21 Mar 2019 22:09:24 +0000 (15:09 -0700)] 
Add support for gang local storage allocation in shared memory

2018-12-11  Julian Brown  <julian@codesourcery.com>
         Chung-Lin Tang  <cltang@codesourcery.com>

     gcc/
     * config/nvptx/nvptx.c (tree-hash-traits.h): Include.
     (gangprivate_shared_size): New global variable.
     (gangprivate_shared_align): Likewise.
     (gangprivate_shared_sym): Likewise.
     (gangprivate_shared_hmap): Likewise.
     (nvptx_option_override): Initialize gangprivate_shared_sym,
     gangprivate_shared_align.
     (nvptx_file_end): Output gangprivate_shared_sym.
     (nvptx_goacc_expand_accel_var): New function.
     (nvptx_set_current_function): New function.
     (TARGET_SET_CURRENT_FUNCTION): Define hook.
     (TARGET_GOACC_EXPAND_ACCEL): Likewise.
     * doc/tm.texi (TARGET_GOACC_EXPAND_ACCEL_VAR): Document new hook.
     * doc/tm.texi.in (TARGET_GOACC_EXPAND_ACCEL_VAR): Likewise.
     * expr.c (expand_expr_real_1): Remap decls marked with the
     "oacc gangprivate" attribute.
     * omp-low.c (omp_context): Add oacc_partitioning_level and
     oacc_addressable_var_decls fields.
     (new_omp_context): Initialize oacc_addressable_var_decls in new
     omp_context.
     (delete_omp_context): Delete oacc_addressable_var_decls in old
     omp_context.
     (lower_oacc_head_tail): Record partitioning-level count in omp context.
     (oacc_record_private_var_clauses, oacc_record_vars_in_bind)
     (mark_oacc_gangprivate): New functions.
     (lower_omp_for): Call oacc_record_private_var_clauses with "for"
     clauses.  Call mark_oacc_gangprivate for gang-partitioned loops.
     (lower_omp_target): Call oacc_record_private_var_clauses with "target"
     clauses.
     Call mark_oacc_gangprivate for offloaded target regions.
     (lower_omp_1): Call vars_in_bind for GIMPLE_BIND within OMP regions.
     * target.def (expand_accel_var): New hook.

     libgomp/
     * testsuite/libgomp.oacc-c-c++-common/gang-private-1.c: New test.
     * testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c: New test.
     * testsuite/libgomp.oacc-c/pr85465.c: New test.
     * testsuite/libgomp.oacc-fortran/gangprivate-attrib-1.f90: New test.kk

5 years agoTweak target selector for libgomp.oacc-c-c++-common/lib-93.c.
Julian Brown [Tue, 19 Mar 2019 13:53:56 +0000 (06:53 -0700)] 
Tweak target selector for libgomp.oacc-c-c++-common/lib-93.c.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/lib-93.c: Adjust target selector.

5 years agoFix hang when running oacc exec with CUDA 9.0 nvprof
Julian Brown [Wed, 27 Feb 2019 00:00:54 +0000 (16:00 -0800)] 
Fix hang when running oacc exec with CUDA 9.0 nvprof

2018-09-20  Tom de Vries  <tdevries@suse.de>
    Cesar Philippidis  <cesar@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.

5 years agoEnable firstprivate OpenACC reductions
Julian Brown [Tue, 26 Feb 2019 23:59:03 +0000 (15:59 -0800)] 
Enable firstprivate OpenACC reductions

2018-09-05  Cesar Philippidis  <cesar@codesourcery.com>
    Chung-Lin Tang  <cltang@codesourcery.com>

gcc/
* gimplify.c (omp_add_variable): Enable firstprivate reduction
variables.

gcc/testsuite/
* c-c++-common/goacc/reduction-8.c: New test.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New
test.
* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New
test.

5 years agoDon't mark OpenACC auto loops as independent inside acc parallel regions
Julian Brown [Tue, 26 Feb 2019 23:55:23 +0000 (15:55 -0800)] 
Don't mark OpenACC auto loops as independent inside acc parallel regions

2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>

gcc/
* omp-low.c (lower_oacc_head_mark): Don't mark OpenACC auto
loops as independent inside acc parallel regions.

gcc/testsuite/
* c-c++-common/goacc/loop-auto-1.c: Adjust test case to conform to
the new behavior of the auto clause in OpenACC 2.5.
* c-c++-common/goacc/loop-auto-2.c: Likewise.
* gcc.dg/goacc/loop-processing-1.c: Likewise.
* c-c++-common/goacc/loop-auto-3.c: New test.
* gfortran.dg/goacc/loop-auto-1.f90: New test.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case
to conform to the new behavior of the auto clause in OpenACC 2.5.

5 years agoEnable GOMP_MAP_FIRSTPRIVATE_INT for OpenACC
Julian Brown [Tue, 26 Feb 2019 23:10:21 +0000 (15:10 -0800)] 
Enable GOMP_MAP_FIRSTPRIVATE_INT for OpenACC

2018-12-22  Cesar Philippidis  <cesar@codesourcery.com>
            Julian Brown  <julian@codesourcery.com>
            Tobias Burnus  <tobias@codesourcery.com>

gcc/
* omp-low.c (maybe_lookup_field_in_outer_ctx): New function.
(convert_to_firstprivate_int): New function.
(convert_from_firstprivate_int): New function.
(lower_omp_target): Enable GOMP_MAP_FIRSTPRIVATE_INT in OpenACC.
Remove unused variable.

libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Handle
GOMP_MAP_FIRSTPRIVATE_INT host addresses.
* plugin/plugin-nvptx.c (nvptx_exec): Handle
GOMP_MAP_FIRSTPRIVATE_INT host addresses.
* testsuite/libgomp.oacc-c++/firstprivate-int.C: New test.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c: New
test.
* testsuite/libgomp.oacc-fortran/firstprivate-int.f90: New test.

5 years agoFix implicit mapping for array slices on lexically-enclosing data constructs (PR70828)
Julian Brown [Tue, 26 Feb 2019 22:22:41 +0000 (14:22 -0800)] 
Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828)

2018-08-28  Julian Brown  <julian@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>

gcc/
* gimplify.c (oacc_array_mapping_info): New struct.
(gimplify_omp_ctx): Add decl_data_clause hash map.
(new_omp_context): Zero-initialise above.
(delete_omp_context): Delete above if allocated.
(gimplify_scan_omp_clauses): Scan for array mappings on data constructs,
and record in above map.
(gomp_oacc_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array
slices) declared in lexically-enclosing data constructs.
* omp-low.c (lower_omp_target): Allow decl for bias not to be present
in OpenACC context.

gcc/testsuite/
* c-c++-common/goacc/acc-data-chain.c: New test.
* gfortran.dg/goacc/pr70828.f90: New test.
* gfortran.dg/goacc/pr70828-2.f90: New test.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.

5 years agoDefault compute dimensions (compile time)
Julian Brown [Tue, 26 Feb 2019 22:12:06 +0000 (14:12 -0800)] 
Default compute dimensions (compile time)

Typo fix relative to last posted version.

2018-10-05  Nathan Sidwell  <nathan@acm.org>
    Tom de Vries  <tdevries@suse.de>
    Thomas Schwinge  <thomas@codesourcery.com>
    Julian Brown  <julian@codesourcery.com>

gcc/
* doc/invoke.texi (fopenacc-dim): Update.
* omp-offload.c (oacc_parse_default_dims): Update.
(oacc_validate_dims): Emit warnings about strange partitioning choices.

gcc/testsuite/
* c-c++-common/goacc/acc-icf.c: Update.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* c-c++-common/goacc/parallel-reduction.c: Likewise.
* c-c++-common/goacc/pr70688.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Add -w.
* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.

5 years agoGenerate sequential loop for OpenACC loop directive inside kernels
Julian Brown [Tue, 26 Feb 2019 21:39:03 +0000 (13:39 -0800)] 
Generate sequential loop for OpenACC loop directive inside kernels

2019-09-20  Chung-Lin Tang <cltang@codesourcery.com>
    Cesar Philippidis  <cesar@codesourcery.com>

gcc/
* omp-expand.c (struct omp_region): Add inside_kernels_p field.
(expand_omp_for_generic): Adjust to generate a 'sequential' loop
when GOMP builtin arguments are BUILT_IN_NONE.
(expand_omp_for): Use expand_omp_for_generic to generate a
non-parallelized loop for OMP_FORs inside OpenACC kernels regions.
(expand_omp): Mark inside_kernels_p field true for regions
nested inside OpenACC kernels constructs.
gcc/testsuite/
* c-c++-common/goacc/kernels-loop-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-2-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-3-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-n-acc-loop.c: New test.
* c-c++-common/goacc/kernels-acc-loop-reduction.c: New test.
* c-c++-common/goacc/kernels-acc-loop-smaller-equal.c: New test.

5 years agoReinstate kernels-restrict behaviour
Julian Brown [Wed, 6 Mar 2019 22:44:56 +0000 (14:44 -0800)] 
Reinstate kernels-restrict behaviour

This patch contains a small fix for upstream churn relative to the last version
posted.

2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
    Julian Brown  <julian@codesourcery.com>

* omp-low.c (install_var_field): New base_pointer_restrict
argument.
(scan_sharing_clauses): Update call to install_var_field.
(omp_target_base_pointers_restrict_p): New function.
(scan_omp_target): Update call to install_var_field.

5 years agoAdjustments and additions to testcases
Julian Brown [Tue, 26 Feb 2019 21:18:36 +0000 (13:18 -0800)] 
Adjustments and additions to testcases

Some additions of redundant "present" clauses dropped.

2018-10-22  Cesar Philippidis  <cesar@codesourcery.com>

gcc/testsuite/
* g++.dg/goacc/loop-1.c: New test.
* g++.dg/goacc/loop-2.c: New test.
* g++.dg/goacc/loop-3.c: New test.

2018-10-22  James Norris  <jnorris@codesourcery.com>
    Cesar Philippidis  <cesar@codesourcery.com>
    Tom de Vries  <tom@codesourcery.com>

libgomp/
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Update parallel
regions to denote variables copyied in via acc enter data as
present.
* testsuite/libgomp.oacc-fortran/data-3.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-4.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/subr.h: Reimplement.
* testsuite/libgomp.oacc-c-c++-common/subr.ptx: Regenerated PTX.
* testsuite/libgomp.oacc-c-c++-common/timer.h: Removed.
* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-71.c: Rework kernel i/f.
* testsuite/libgomp.oacc-c-c++-common/lib-72.c: Rework kernel i/f and
change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-74.c: Rework kernel i/f and
timing checks.
* testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-77.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-80.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-93.c: New test.

5 years agoVarious OpenACC reduction enhancements - test cases
Julian Brown [Tue, 12 Feb 2019 23:14:22 +0000 (15:14 -0800)] 
Various OpenACC reduction enhancements - test cases

2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
            Nathan Sidwell  <nathan@acm.org>
            Julian Brown  <julian@codesourcery.com>

        gcc/testsuite/
        * c-c++-common/goacc/orphan-reductions-1.c: New test.
        * c-c++-common/goacc/reduction-7.c: New test.
        * c-c++-common/goacc/routine-4.c: Update.
        * g++.dg/goacc/reductions-1.C: New test.
        * gcc.dg/goacc/loop-processing-1.c: Update.
        * gfortran.dg/goacc/orphan-reductions-1.f90: New test.

        libgomp/
        * libgomp.oacc-c-c++-common/par-reduction-3.c: New test.
        * libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test.
        * libgomp.oacc-fortran/reduction-9.f90: New test.

5 years agoVarious OpenACC reduction enhancements - ME and nvptx changes
Julian Brown [Tue, 12 Feb 2019 23:06:55 +0000 (15:06 -0800)] 
Various OpenACC reduction enhancements - ME and nvptx changes

Parts of the first posting got lost in the second posting, above.
This version hopefully contains everything.

2018-10-30  Cesar Philippidis  <cesar@codesourcery.com>

gcc/
* config/nvptx/nvptx.c (nvptx_propagate_unified): New.
(nvptx_split_blocks): Call it for cond_uni insn.
(nvptx_expand_cond_uni): New.
(enum nvptx_builtins): Add NVPTX_BUILTIN_COND_UNI.
(nvptx_init_builtins): Initialize it.
(nvptx_expand_builtin):
(nvptx_generate_vector_shuffle): Change integral SHIFT operand to
tree BITS operand.
(nvptx_vector_reduction): New.
(nvptx_adjust_reduction_type): New.
(nvptx_goacc_reduction_setup): Use it to adjust the type of ref_to_res.
(nvptx_goacc_reduction_init): Don't update LHS if it doesn't exist.
(nvptx_goacc_reduction_fini): Call nvptx_vector_reduction for vector.
Use it to adjust the type of ref_to_res.
(nvptx_goacc_reduction_teardown):
* config/nvptx/nvptx.md (cond_uni): New pattern.

5 years agoVarious OpenACC reduction enhancements - FE changes
Julian Brown [Tue, 12 Feb 2019 22:56:12 +0000 (14:56 -0800)] 
Various OpenACC reduction enhancements - FE changes

This version differs somewhat from the last version posted upstream
(and addresses some of Jakub's review comments).

2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
    Nathan Sidwell  <nathan@acm.org>
    Julian Brown  <julian@codesourcery.com>

gcc/c/
* c-parser.c (c_parser_omp_variable_list): New c_omp_region_type
argument.  Use it to specialize handling of OMP_CLAUSE_REDUCTION for
OpenACC.
(c_parser_oacc_data_clause): Add region-type argument.
(c_parser_oacc_data_clause_deviceptr): Likewise.
(c_parser_omp_clause_reduction): Change is_omp boolean parameter to
c_omp_region_type.  Update call to c_parser_omp_variable_list.
(c_parser_oacc_all_clauses): Update calls to
c_parser_omp_clause_reduction.
(c_parser_omp_all_clauses): Likewise.
(c_parser_oacc_cache): Update call to c_parser_omp_var_list_parens.
* c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC
gang reductions.  Suppress user-defined reduction error for OpenACC.

gcc/cp/
* parser.c (cp_parser_omp_var_list_no_open):  New c_omp_region_type
argument.  Use it to specialize handling of OMP_CLAUSE_REDUCTION for
OpenACC.
(cp_parser_omp_var_list): Add c_omp_region_type argument. Update call
to cp_parser_omp_var_list_parens.
(cp_parser_oacc_data_clause): Update call to cp_parser_omp_var_list.
(cp_parser_omp_clause_reduction): Change is_omp boolean parameter to
c_omp_region_type.  Update call to cp_parser_omp_var_list_no_open.
(cp_parser_oacc_all_clauses): Update call to
cp_parser_omp_clause_reduction.
(cp_parser_omp_all_clauses): Likewise.
* semantics.c (finish_omp_reduction_clause): Add c_omp_region_type
argument.  Suppress user-defined reduction error for OpenACC.
(finish_omp_clauses): Emit an error on orphan OpenACC gang reductions.

gcc/fortran/
* openmp.c (resolve_oacc_loop_blocks): Emit an error on orphan OpenACC
gang reductions.
* trans-openmp.c (gfc_omp_clause_copy_ctor): Permit reductions.

5 years agoAdd OpenACC Fortran support for deviceptr and variable in common blocks
Julian Brown [Tue, 12 Feb 2019 22:32:34 +0000 (14:32 -0800)] 
Add OpenACC Fortran support for deviceptr and variable in common blocks

2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
    James Norris  <jnorris@codesourcery.com>

gcc/fortran/
* openmp.c (gfc_match_omp_map_clause): Re-write handling of the
deviceptr clause.  Add new common_blocks argument.  Propagate it to
gfc_match_omp_variable_list.
(gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses.
(resolve_positive_int_expr): Promote the warning to an error.
(check_array_not_assumed): Remove pointer check.
(resolve_oacc_nested_loops): Error on do concurrent loops.
* trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data
mappings for deviceptr clauses.
(gfc_trans_omp_clauses): Likewise.

gcc/
* gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
(oacc_default_clause): Privatize fortran common blocks.
(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
Defer the expansion of DECL_VALUE_EXPR for common block decls.
(gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
appropriate.
(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
implicit deviceptr mappings.

gcc/testsuite/
* c-c++-common/goacc/deviceptr-4.c: Update.
* gfortran.dg/goacc/common-block-1.f90: New test.
* gfortran.dg/goacc/common-block-2.f90: New test.
* gfortran.dg/goacc/loop-2.f95: Update.
* gfortran.dg/goacc/loop-3-2.f95: Update.
* gfortran.dg/goacc/loop-3.f95: Update.
* gfortran.dg/goacc/loop-5.f95: Update.
* gfortran.dg/goacc/pr72715.f90: New test.
* gfortran.dg/goacc/sie.f95: Update.
* gfortran.dg/goacc/tile-1.f90: Update.
* gfortran.dg/gomp/pr77516.f90: Update.

libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr
clause.
(GOACC_data_start): Likewise.
* testsuite/libgomp.oacc-fortran/common-block-1.f90: New test.
* testsuite/libgomp.oacc-fortran/common-block-2.f90: New test.
* testsuite/libgomp.oacc-fortran/common-block-3.f90: New test.
* testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test.

5 years agoTweak error return value for acc_set_cuda_stream.
Julian Brown [Tue, 12 Feb 2019 14:36:03 +0000 (06:36 -0800)] 
Tweak error return value for acc_set_cuda_stream.

The return value of acc_set_cuda_stream is unspecified in OpenACC 2.6.
The testsuite changes might be unnecessary with the current async code.

libgomp/
* oacc-cuda.c (acc_set_cuda_stream): Return 0 on error/invalid
arguments.
* testsuite/libgomp.oacc-c-c++-common/lib-84.c: Handle unnumbered
async stream being an alias for a numbered async stream.
* testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise.

5 years agoAdd support for OpenACC routine nohost clause
Julian Brown [Thu, 7 Feb 2019 01:46:25 +0000 (17:46 -0800)] 
Add support for OpenACC routine nohost clause

2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
    Cesar Philippidis  <cesar@codesourcery.com>

gcc/
* tree-core.h (omp_clause_code): Add OMP_CLAUSE_NOHOST.
* tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
Update for these.
* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_NOHOST.
* gimplify.c (gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOHOST.
* tree-nested.c (convert_nonlocal_omp_clauses)
(convert_local_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Likewise.
* omp-offload.c (maybe_discard_oacc_function): New function.
(execute_oacc_device_lower) [!ACCEL_COMPILER]: Handle OpenACC
nohost clauses.

gcc/c-family/
* c-attribs.c (c_common_attribute_table): Set min_len to -1 for
"omp declare target".
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_NOHST.

gcc/c/
* c-parser.c (c_parser_omp_clause_name): Handle "nohost".
(c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST.
(c_parser_oacc_routine, c_finish_oacc_routine): Update.
* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_NOHOST.

gcc/cp/
* parser.c (cp_parser_omp_clause_name): Handle "nohost".
(cp_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST,
(cp_parser_oacc_routine, cp_finalize_oacc_routine): Update.
* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NOHOST.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_NOHOST.

gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add nohost members.
* openmp.c (omp_mask2): Add OMP_CLAUSE_NOHOST.
(gfc_match_omp_clauses): Handle OMP_CLAUSE_NOHOST.
(gfc_match_oacc_routine): Set oacc_function_nohost when appropriate.
* gfortran.h (symbol_attribute): Add oacc_function_nohost member.
* trans-openmp.c (gfc_add_omp_offload_attributes): Use it to decide
whether to generate an OMP_CLAUSE_NOHOST clause.
(gfc_trans_omp_clauses_1): Unreachable code to generate an
OMP_CLAUSE_NOHOST clause.

gcc/testsuite/
* c-c++-common/goacc/classify-routine.c: Adjust test.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-2.c: Likewise.
* c-c++-common/goacc/routine-nohost-1.c: New test.
* g++.dg/goacc/routine-2.C: Adjust test.
* gfortran.dg/goacc/pr72741.f90: New test.

libgomp/
* testsuite/libgomp.oacc-c-c++-common/routine-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c:
Update test.
* testsuite/libgomp.oacc-fortran/routine-6.f90: Likewise.

5 years agoMerge non-contiguous array support patches.
Chung-Lin Tang [Sun, 19 Apr 2020 12:10:43 +0000 (05:10 -0700)] 
Merge non-contiguous array support patches.

This version is based from v4, posted upstream here:
https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html

2020-04-19  Chung-Lin Tang  <cltang@codesourcery.com>

PR other/76739

gcc/c/
* c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
parameter, adjust recursive call site, add cases for allowing
pointer based multi-dimensional arrays for OpenACC.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.

gcc/cp/
* semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
parameter, adjust recursive call site, add cases for allowing
pointer based multi-dimensional arrays for OpenACC.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.

gcc/fortran/
* f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol.
* types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.

gcc/
* builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
* omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
* gimplify.c (gimplify_scan_omp_clauses): Skip gimplification of
OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
* omp-expand.c (expand_omp_target): Add non-contiguous array descriptor
pointers to variadic arguments.
* omp-low.c (append_field_to_record_type): New function.
(create_noncontig_array_descr_type): Likewise.
(create_noncontig_array_descr_init_code): Likewise.
(scan_sharing_clauses): For non-contiguous array map kinds, check for
supported dimension structure, and install non-contiguous array
variable into current omp_context.
(reorder_noncontig_array_clauses): New function.
(scan_omp_target): Call reorder_noncontig_array_clauses to place
non-contiguous array map clauses at beginning of clause sequence.
(lower_omp_target): Add handling for non-contiguous array map kinds,
add all created non-contiguous array descriptors to
gimple_omp_target_data_arg.

gcc/testsuite/
* c-c++-common/goacc/noncontig_array-1.c: New test.

libgomp/
* libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration.
* libgomp.h (gomp_map_vars_openacc): New function declaration.
* oacc-int.h (struct goacc_ncarray_dim): New struct declaration.
(struct goacc_ncarray_descr_type): Likewise.
(struct goacc_ncarray): Likewise.
(struct goacc_ncarray_info): Likewise.
(goacc_noncontig_array_create_ptrblock): New function declaration.
* oacc-parallel.c (goacc_noncontig_array_count_rows): New function.
(goacc_noncontig_array_compute_sizes): Likewise.
(goacc_noncontig_array_fill_rows_1): Likewise.
(goacc_noncontig_array_fill_rows): Likewise.
(goacc_process_noncontiguous_arrays): Likewise.
(goacc_noncontig_array_create_ptrblock): Likewise.
(GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to
handle non-contiguous array descriptors at end of varargs, adjust
to use gomp_map_vars_openacc.
(GOACC_data_start): Likewise. Adjust function type to accept varargs.
* target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info *
nca_info parameter, add handling code for non-contiguous arrays.
(gomp_map_vars_openacc): Add new function for specialization of
gomp_map_vars_internal for OpenACC structured region usage.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support
header for new tests.

include/
* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define.
(enum gomp_map_kind): Add GOMP_MAP_NONCONTIG_ARRAY,
GOMP_MAP_NONCONTIG_ARRAY_TO, GOMP_MAP_NONCONTIG_ARRAY_FROM,
GOMP_MAP_NONCONTIG_ARRAY_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO,
GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM,
GOMP_MAP_NONCONTIG_ARRAY_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC,
GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT.
(GOMP_MAP_NONCONTIG_ARRAY_P): Define.

5 years agoOpenACC reference count consistency checking
Julian Brown [Fri, 30 Nov 2018 10:40:49 +0000 (02:40 -0800)] 
OpenACC reference count consistency checking

2018-11-28  Julian Brown  <julian@codesourcery.com>
    Tobias Burnus  <tobias@codesourcery.com>

libgomp/
* libgomp.h (RC_CHECKING): New macro, disabled by default, guarding all
hunks in this patch.
(target_mem_desc): Add forward declaration.
(async_tgt_use): New struct.
(target_mem_desc): Add refcount_chk, mark fields.
(acc_dispatch_t): Add tgt_uses, au_lock fields.
(dump_tgt, gomp_rc_check): Add prototypes.
* oacc-async (goacc_async_unmap_tgt): Add refcount self-check code.
(goacc_async_copyout_unmap_vars): Likewise.
(goacc_remove_var_async): Likewise.
* oacc-parallel.c (GOACC_parallel_keyed_internal): Add refcount
self-check code.
(GOACC_data_start, GOACC_data_end, GOACC_enter_exit_data): Likewise.
* target.c (stdio.h): Include.
(dump_tgt, rc_check_clear, rc_check_count, rc_check_verify)
(gomp_rc_check): New functions to consistency-check reference counts.
(gomp_target_init): Initialise self-check-related device fields.

5 years agoc++ ICE with nested requirement as default tpl parm[PR94827]
Nathan Sidwell [Thu, 30 Apr 2020 15:23:16 +0000 (08:23 -0700)] 
c++ ICE with nested requirement as default tpl parm[PR94827]

Template headers are not incrementally updated as we parse its parameters.
We maintain a dummy level until the closing > when we replace the dummy with
a real parameter set.  requires processing was expecting a properly populated
arg_vec in current_template_parms, and then creates a self-mapping of parameters
from that.  But we don't need to do that, just teach map_arguments to look at
TREE_VALUE when args is NULL.

* constraint.cc (map_arguments): If ARGS is null, it's a
self-mapping of parms.
(finish_nested_requirement): Do not pass argified
current_template_parms to normalization.
(tsubst_nested_requirement): Don't assert no template parms.

5 years agolibstdc++: Avoid errors in allocator's noexcept-specifier (PR 89510)
Jonathan Wakely [Thu, 30 Apr 2020 14:47:52 +0000 (15:47 +0100)] 
libstdc++: Avoid errors in allocator's noexcept-specifier (PR 89510)

This fixes a regression due to the conditional noexcept-specifier on
std::allocator::construct and std::allocator::destroy, as well as the
corresponding members of new_allocator, malloc_allocator, and
allocator_traits. Those noexcept-specifiers were using expressions which
might be ill-formed, which caused errors outside the immediate context
when checking for the presence of construct and destroy in SFINAE
contexts.

The fix is to use the is_nothrow_constructible and
is_nothrow_destructible type traits instead, because those traits are
safe to use even when the construction/destruction itself is not valid.

The is_nothrow_constructible trait will be false for a type that is not
also nothrow-destructible, even if the new-expression used in the
construct function body is actually noexcept. That's not the correct
answer, but isn't a problem because providing a noexcept-specifier on
these functions is not required by the standard anyway. If the answer is
false when it should be true, that's suboptimal but OK (unlike giving
errors for valid code, or giving a true answer when it should be false).

PR libstdc++/89510
* include/bits/alloc_traits.h (allocator_traits::_S_construct)
(allocator_traits::_S_destroy)
(allocator_traits<allocator<T>>::construct): Use traits in
noexcept-specifiers.
* include/bits/allocator.h (allocator<void>::construct)
(allocator<void>::destroy): Likewise.
* include/ext/malloc_allocator.h (malloc_allocator::construct)
(malloc_allocator::destroy): Likewise.
* include/ext/new_allocator.h (new_allocator::construct)
(new_allocator::destroy): Likewise.
* testsuite/20_util/allocator/89510.cc: New test.
* testsuite/ext/malloc_allocator/89510.cc: New test.
* testsuite/ext/new_allocator/89510.cc: New test.

5 years agocoroutines: Fix handling of artificial vars [PR94886]
Iain Sandoe [Thu, 30 Apr 2020 09:42:36 +0000 (10:42 +0100)] 
coroutines: Fix handling of artificial vars [PR94886]

The testcase ICEs because the range-based for generates three
artificial variables that need to be allocated to the coroutine
frame but, when walking the BIND_EXR that contains these, the
DECL_INITIAL for one of them refers to an entry appearing later,
which means that the frame entry hasn't been allocated when that
INITIAL is walked.

The solution is to defer walking the DECL_INITIAL/SIZE etc. until
all the BIND_EXPR vars have been processed.

gcc/cp/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94886
* coroutines.cc (transform_local_var_uses): Defer walking
the DECL_INITIALs of BIND_EXPR vars until all the frame
allocations have been made.

gcc/testsuite/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94886
* g++.dg/coroutines/pr94886-folly-3.C: New test.

5 years agocoroutines: Fix handling of target cleanup exprs [PR94883]
Iain Sandoe [Wed, 29 Apr 2020 18:46:35 +0000 (19:46 +0100)] 
coroutines: Fix handling of target cleanup exprs [PR94883]

The problem here is that target cleanup expressions have been
added to the initialisers for the awaitable (and returns of
non-trivial values from await_suspend() calls.  This is because
the expansion of the co_await into its control flow is not
apparent to the machinery adding the target cleanup expressions.
The solution being tested is simply to recreate target expressions
as the co_awaits are lowered.  Teaching the machinery to handle
walking co_await expressions in different ways at different points
(outside the coroutine transformation) seems overly complex.

gcc/cp/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94883
* coroutines.cc (register_awaits): Update target
expressions for awaitable and suspend handle
initializers.

gcc/testsuite/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94883
* g++.dg/coroutines/pr94883-folly-2.C: New test.

5 years agocoroutines: Fix cases where proxy variables are used [PR94879]
Iain Sandoe [Mon, 27 Apr 2020 14:21:25 +0000 (15:21 +0100)] 
coroutines: Fix cases where proxy variables are used [PR94879]

There are several places where the handling of a variable
declaration depends on whether it corresponds to a compiler
temporary, or to some other entity.  We were testing that var
decls were artificial in determining this.  However, proxy vars
are also artificial so that this is not sufficient.  The solution
is to exclude variables with a DECL_VALUE_EXPR as well, since
the value variable will not be a temporary.

gcc/cp/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94879
* coroutines.cc (build_co_await): Account for variables
with DECL_VALUE_EXPRs.
(captures_temporary): Likewise.
(register_awaits): Likewise.

gcc/testsuite/ChangeLog:

2020-04-30  Iain Sandoe  <iain@sandoe.co.uk>

PR c++/94879
* g++.dg/coroutines/pr94879-folly-1.C: New test.

5 years agodiagnostics: Fix spelling in comment
Jonathan Wakely [Thu, 30 Apr 2020 13:42:24 +0000 (14:42 +0100)] 
diagnostics: Fix spelling in comment

gcc/ChangeLog:

* pretty-print.c (pp_take_prefix): Fix spelling in comment.

5 years agotree: Don't reuse types if TYPE_USER_ALIGN differ [PR94775]
Marek Polacek [Wed, 29 Apr 2020 19:36:35 +0000 (15:36 -0400)] 
tree: Don't reuse types if TYPE_USER_ALIGN differ [PR94775]

Here we trip on the TYPE_USER_ALIGN (t) assert in strip_typedefs: it
gets "const d[0]" with TYPE_USER_ALIGN=0 but the result built by
build_cplus_array_type is "const char[0]" with TYPE_USER_ALIGN=1.

When we strip_typedefs the element of the array "const d", we see it's
a typedef_variant_p, so we look at its DECL_ORIGINAL_TYPE, which is
char, but we need to add the const qualifier, so we call
cp_build_qualified_type -> build_qualified_type
where get_qualified_type checks to see if we already have such a type
by walking the variants list, which in this case is:

  char -> c -> const char -> const char -> d -> const d

Because check_base_type only checks TYPE_ALIGN and not TYPE_USER_ALIGN,
we choose the first const char, which has TYPE_USER_ALIGN set.  If the
element type of an array has TYPE_USER_ALIGN, the array type gets it too.

So we can make check_base_type stricter.  I was afraid that it might make
us reuse types less often, but measuring showed that we build the same
amount of types with and without the patch, while bootstrapping.

PR c++/94775
* tree.c (check_base_type): Return true only if TYPE_USER_ALIGN match.
(check_aligned_type): Check if TYPE_USER_ALIGN match.

* g++.dg/warn/Warray-bounds-10.C: New test.

5 years ago[AArch64] Make -moutline-atomics on by default
Kyrylo Tkachov [Thu, 30 Apr 2020 12:12:13 +0000 (13:12 +0100)] 
[AArch64] Make -moutline-atomics on by default

2020-04-30  Kyrylo Tkachov  <kyrylo.tkachov@arm.com>

* config/aarch64/aarch64.h (TARGET_OUTLINE_ATOMICS): Define.
* config/aarch64/aarch64.opt (moutline-atomics): Change to Int variable.
* doc/invoke.texi (moutline-atomics): Document as on by default.

5 years agoaarch64: don't emit bti j after NOTE_INSN_DELETED_LABEL [PR94748]
Szabolcs Nagy [Fri, 24 Apr 2020 16:36:02 +0000 (17:36 +0100)] 
aarch64: don't emit bti j after NOTE_INSN_DELETED_LABEL [PR94748]

It was previously discussed that indirect branches cannot go to
NOTE_INSN_DELETED_LABEL so inserting a landing pad is unnecessary.
See https://gcc.gnu.org/pipermail/gcc-patches/2019-May/522625.html

Before the patch a bti j was inserted after the label in

  __attribute__((target("branch-protection=bti")))
  int foo (void)
  {
  label:
    return 0;
  }

This is not necessary and weakens the security protection.

gcc/ChangeLog:

PR target/94748
* config/aarch64/aarch64-bti-insert.c (rest_of_insert_bti): Remove
the check for NOTE_INSN_DELETED_LABEL.

gcc/testsuite/ChangeLog:

PR target/94748
* gcc.target/aarch64/pr94748.c: New test.

5 years agod: Fix documentation of -defaultlib= and -debuglib=
Iain Buclaw [Thu, 30 Apr 2020 10:21:09 +0000 (12:21 +0200)] 
d: Fix documentation of -defaultlib= and -debuglib=

From the generated manpages, it was not clear that its usage is
'-debuglib=<libname>'.

gcc/d/ChangeLog:

* gdc.texi (Options for Linking): Clarify usage of -defaultlib= and
-debuglib= options.

5 years agoAdd missing ChangeLog entry
Iain Buclaw [Thu, 30 Apr 2020 10:20:04 +0000 (12:20 +0200)] 
Add missing ChangeLog entry

5 years agod: Merge upstream dmd 934df6f8c, druntime 7bdd83d7
Iain Buclaw [Thu, 30 Apr 2020 09:03:33 +0000 (11:03 +0200)] 
d: Merge upstream dmd 934df6f8c, druntime 7bdd83d7

Corrects a previous change made to the SPARC stdc bindings, and
backports PPC-related fixes.  The library and language testsuite now
passes fully on powerpc64le-linux-gnu.

Fixes: PR d/90719
Fixes: PR d/94825
Reviewed-on: https://github.com/dlang/dmd/pull/11079
     https://github.com/dlang/druntime/pull/3078
     https://github.com/dlang/druntime/pull/3083

libphobos/ChangeLog:

PR d/94825
* libdruntime/Makefile.am (DRUNTIME_SOURCES_CONFIGURED): Remove
config/powerpc/switchcontext.S
* libdruntime/Makefile.in: Regenerate.
* libdruntime/config/powerpc/callwithstack.S: Remove.
* libdruntime/config/powerpc/switchcontext.S: Fix symbol name of
fiber_switchContext.
* libdruntime/core/thread.d: Disable fiber migration tests on PPC.
* testsuite/libphobos.thread/fiber_guard_page.d: Set guardPageSize
same as stackSize.

5 years ago--with-{documentation,changes}-root-url tweaks
Jakub Jelinek [Thu, 30 Apr 2020 09:49:40 +0000 (11:49 +0200)] 
--with-{documentation,changes}-root-url tweaks

>   , CHANGES_URL ("gcc-10/changes.html#empty_base");
>
> where the macro would just use preprocessor string concatenation?

Ok, the following patch implements it (doesn't introduce a separate
macro and just uses CHANGES_ROOT_URL "gcc-10/changes.html#empty_base"),
in addition adds the documentation Joseph requested.

2020-04-30  Jakub Jelinek  <jakub@redhat.com>

* configure.ac (--with-documentation-root-url,
--with-changes-root-url): Diagnose URL not ending with /,
use AC_DEFINE_UNQUOTED instead of AC_SUBST.
* opts.h (get_changes_url): Remove.
* opts.c (get_changes_url): Remove.
* Makefile.in (CFLAGS-opts.o): Don't add -DDOCUMENTATION_ROOT_URL
or -DCHANGES_ROOT_URL.
* doc/install.texi (--with-documentation-root-url,
--with-changes-root-url): Document.
* config/arm/arm.c (aapcs_vfp_is_call_or_return_candidate): Don't call
get_changes_url and free, change url variable type to const char * and
set it to CHANGES_ROOT_URL "gcc-10/changes.html#empty_base".
* config/s390/s390.c (s390_function_arg_vector,
s390_function_arg_float): Likewise.
* config/aarch64/aarch64.c (aarch64_vfp_is_call_or_return_candidate):
Likewise.
* config/rs6000/rs6000-call.c (rs6000_discover_homogeneous_aggregate):
Likewise.
* config.in: Regenerate.
* configure: Regenerate.

5 years agoPR94725 - deleting gfortran.dg/dtio_5.f90 and pdt_5.f03
Paul Thomas [Thu, 30 Apr 2020 08:56:01 +0000 (09:56 +0100)] 
PR94725 - deleting gfortran.dg/dtio_5.f90 and pdt_5.f03

5 years agoarm: Remove duplicate entries in isr_attribute_args [PR target/57002]
Christophe Lyon [Wed, 29 Apr 2020 09:55:45 +0000 (09:55 +0000)] 
arm: Remove duplicate entries in isr_attribute_args [PR target/57002]

Remove two duplicate entries in isr_attribute_args ("abort" and
"ABORT").

2020-04-30  Christophe Lyon  <christophe.lyon@linaro.org>

PR target/57002
gcc/
* config/arm/arm.c (isr_attribute_args): Remove duplicate entries.

5 years agoIBM Z: vec_store_len_r/vec_load_len_r fix
Andreas Krebbel [Thu, 30 Apr 2020 06:29:26 +0000 (08:29 +0200)] 
IBM Z: vec_store_len_r/vec_load_len_r fix

This fixes a problem with the vec_store_len_r intrinsic.  The macros
mapping the intrinsic to a GCC builtin had the wrong signature.

With the patch an immediate length operand of vlrl/vstrl is handled
the same way as if it was passed in a register to vlrlr/vstrlr.
Values bigger than 15 always load the full vector.  If it can be
recognized that it is in effect a full vector register load or store
it is now implemented with vl/vst instead.

gcc/ChangeLog:

2020-04-30  Andreas Krebbel  <krebbel@linux.ibm.com>

* config/s390/constraints.md ("j>f", "jb4"): New constraints.
* config/s390/vecintrin.h (vec_load_len_r, vec_store_len_r): Fix
macro definitions.
* config/s390/vx-builtins.md ("vlrlrv16qi", "vstrlrv16qi"): Add a
separate expander.
("*vlrlrv16qi", "*vstrlrv16qi"): Add alternative for vl/vst.
Change constraint for vlrl/vstrl to jb4.

gcc/testsuite/ChangeLog:

2020-04-30  Andreas Krebbel  <krebbel@linux.ibm.com>

* gcc.target/s390/zvector/vec_load_len_r.c: New test.
* gcc.target/s390/zvector/vec_store_len_r.c: New test.

5 years agovar-tracking.c: Fix possible use of uninitialized variable pre
Stefan Schulze Frielinghaus [Sun, 26 Apr 2020 15:10:51 +0000 (17:10 +0200)] 
var-tracking.c: Fix possible use of uninitialized variable pre

While bootstrapping GCC on S/390 the following warning/error is raised:

gcc/var-tracking.c:10239:34: error: 'pre' may be used uninitialized in this function [-Werror=maybe-uninitialized]
10239 |       VTI (bb)->out.stack_adjust += pre;
      |                                  ^

The lines of interest are:

  HOST_WIDE_INT pre, post = 0;
  // ...
  if (!frame_pointer_needed)
    {
      insn_stack_adjust_offset_pre_post (insn, &pre, &post);
      // ...
    }

  // ...
  adjust_insn (bb, insn);

  if (!frame_pointer_needed && pre)
    VTI (bb)->out.stack_adjust += pre;

Both if statements depend on global variable frame_pointer_needed.  In function
insn_stack_adjust_offset_pre_post local variable pre is initialized.  The
problematic part is the function call between both if statements.  Since
adjust_insn also calls functions which are defined in a different compilation
unit, we are not able to prove that global variable frame_pointer_needed is not
altered by adjust_insn and its siblings.  Thus we must assume that
frame_pointer_needed may be true before the call and false afterwards which
renders the warning true (admitted the location hint is not totally perfect).
By initialising pre we silence the warning.

gcc/ChangeLog:

2020-04-30  Stefan Schulze Frielinghaus  <stefansf@linux.ibm.com>

        * var-tracking.c (vt_initialize): Move variables pre and post
        into inner block and initialize both in order to fix warning
        about uninitialized use.  Remove unnecessary checks for
        frame_pointer_needed.

5 years agotoplev.c: Check for null argument to fprintf
Stefan Schulze Frielinghaus [Sun, 26 Apr 2020 07:26:39 +0000 (09:26 +0200)] 
toplev.c: Check for null argument to fprintf

Ensure that CF does not equal NULL in function output_stack_usage_1
before calling fprintf.  This fixes the following warning/error:

gcc/toplev.c:976:13: error: argument 1 null where non-null expected [-Werror=nonnull]
  976 |     fprintf (cf, "\\n" HOST_WIDE_INT_PRINT_DEC " bytes (%s)",
      |     ~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  977 |       stack_usage,
      |       ~~~~~~~~~~~~
  978 |       stack_usage_kind_str[stack_usage_kind]);

An example call side where CF is NULL is in function output_stack_usage.

gcc/ChangeLog:

2020-04-30  Stefan Schulze Frielinghaus  <stefansf@linux.ibm.com>

* toplev.c (output_stack_usage_1): Ensure that first
argument to fprintf is not null.

5 years agoDaily bump.
GCC Administrator [Thu, 30 Apr 2020 00:16:16 +0000 (00:16 +0000)] 
Daily bump.

5 years agodiagnostics: Add %{...%} pretty-format support for URLs and use it in -Wpsabi diagnostics
Jakub Jelinek [Wed, 29 Apr 2020 20:41:47 +0000 (22:41 +0200)] 
diagnostics: Add %{...%} pretty-format support for URLs and use it in -Wpsabi diagnostics

The following patch attempts to use the diagnostics URL support if available
to provide more information about the C++17 empty base and C++20
[[no_unique_address]] empty class ABI changes in -Wpsabi diagnostics.

in GCC 10.1 at the end of the diagnostics is then in some terminals
underlined with a dotted line and points to a (to be written) anchor in
gcc-10/changes.html which we need to write anyway.

2020-04-29  Jakub Jelinek  <jakub@redhat.com>

* configure.ac (-with-changes-root-url): New configure option,
defaulting to https://gcc.gnu.org/.
* Makefile.in (CFLAGS-opts.o): Define CHANGES_ROOT_URL for
opts.c.
* pretty-print.c (get_end_url_string): New function.
(pp_format): Handle %{ and %} for URLs.
(pp_begin_url): Use pp_string instead of pp_printf.
(pp_end_url): Use get_end_url_string.
* opts.h (get_changes_url): Declare.
* opts.c (get_changes_url): New function.
* config/rs6000/rs6000-call.c: Include opts.h.
(rs6000_discover_homogeneous_aggregate): Use %{in GCC 10.1%} instead
of just in GCC 10.1 in diagnostics and add URL.
* config/arm/arm.c (aapcs_vfp_is_call_or_return_candidate): Likewise.
* config/aarch64/aarch64.c (aarch64_vfp_is_call_or_return_candidate):
Likewise.
* config/s390/s390.c (s390_function_arg_vector,
s390_function_arg_float): Likewise.
* configure: Regenerated.

* c-format.c (PP_FORMAT_CHAR_TABLE): Add %{ and %}.

5 years agos390: Fix up -Wpsabi diagnostics + [[no_unique_address]] empty member fix [PR94704]
Jakub Jelinek [Wed, 29 Apr 2020 20:38:01 +0000 (22:38 +0200)] 
s390: Fix up -Wpsabi diagnostics + [[no_unique_address]] empty member fix [PR94704]

So, based on the yesterday's discussions, similarly to powerpc64le-linux
I've done some testing for s390x-linux too.

First of all, I found a bug in my patch from yesterday, it was printing
the wrong type like 'double' etc. rather than the class that contained such
the element.  Fix below.

For s390x-linux, I was using
struct X { };
struct Y { int : 0; };
struct Z { int : 0; Y y; };
struct U : public X { X q; };
struct A { double a; };
struct B : public X { double a; };
struct C : public Y { double a; };
struct D : public Z { double a; };
struct E : public U { double a; };
struct F { [[no_unique_address]] X x; double a; };
struct G { [[no_unique_address]] Y y; double a; };
struct H { [[no_unique_address]] Z z; double a; };
struct I { [[no_unique_address]] U u; double a; };
struct J { double a; [[no_unique_address]] X x; };
struct K { double a; [[no_unique_address]] Y y; };
struct L { double a; [[no_unique_address]] Z z; };
struct M { double a; [[no_unique_address]] U u; };
 #define T(S, s) extern S s; extern void foo##s (S); int bar##s () { foo##s (s); return 0; }
T (A, a)
T (B, b)
T (C, c)
T (D, d)
T (E, e)
T (F, f)
T (G, g)
T (H, h)
T (I, i)
T (J, j)
T (K, k)
T (L, l)
T (M, m)
as testcase and looking for "\tld\t%f0,".
While g++ 9 with -std=c++17 used to pass in fpr just
A, g++ 9 -std=c++14, as well as current trunk -std=c++14 & 17
and clang++ from today -std=c++14 & 17 all pass A, B, C
in fpr and nothing else.  The intent stated by Jason seems to be
that A, B, C, F, G, J, K should all be passed in fpr.

Attached are two (updated) versions of the patch on top of the
powerpc+middle-end patch just posted.

The first one emits two separate -Wpsabi warnings like powerpc, one for
the -std=c++14 vs. -std=c++17 ABI difference and one for GCC 9 vs. 10
[[no_unique_address]] passing changes, the other one is silent about the
second case.

2020-04-29  Jakub Jelinek  <jakub@redhat.com>

PR target/94704
* config/s390/s390.c (s390_function_arg_vector,
s390_function_arg_float): Use DECL_FIELD_ABI_IGNORED instead of
cxx17_empty_base_field_p.  In -Wpsabi diagnostics use the type
passed to the function rather than the type of the single element.
Rename cxx17_empty_base_seen variable to empty_base_seen, change
type to int, and adjust diagnostics depending on if the field
has [[no_unique_attribute]] or not.

* g++.target/s390/s390.exp: New file.
* g++.target/s390/pr94704-1.C: New test.
* g++.target/s390/pr94704-2.C: New test.
* g++.target/s390/pr94704-3.C: New test.
* g++.target/s390/pr94704-4.C: New test.

5 years agolibstdc++: Fix outdated comment about std::string instantiations (PR 94854)
Jonathan Wakely [Wed, 29 Apr 2020 17:57:34 +0000 (18:57 +0100)] 
libstdc++: Fix outdated comment about std::string instantiations (PR 94854)

PR libstdc++/94854
* include/bits/basic_string.tcc: Update comment about explicit
instantiations.

5 years agox86: Fix -O0 remaining intrinsic macros [PR94832]
Jakub Jelinek [Wed, 29 Apr 2020 15:31:26 +0000 (17:31 +0200)] 
x86: Fix -O0 remaining intrinsic macros [PR94832]

A few other macros seem to suffer from the same issue.  What I've done was:
cat gcc/config/i386/*intrin.h | sed -e ':x /\\$/ { N; s/\\\n//g ; bx }' \
| grep '^[[:blank:]]*#[[:blank:]]*define[[:blank:]].*(' | sed 's/[  ]\+/ /g' \
> /tmp/macros
and then looking for regexps:
)[a-zA-Z]
) [a-zA-Z]
[a-zA-Z][-+*/%]
[a-zA-Z] [-+*/%]
[-+*/%][a-zA-Z]
[-+*/%] [a-zA-Z]
in the resulting file.

2020-04-29  Jakub Jelinek  <jakub@redhat.com>

PR target/94832
* config/i386/avx512bwintrin.h (_mm512_alignr_epi8,
_mm512_mask_alignr_epi8, _mm512_maskz_alignr_epi8): Wrap macro operands
used in casts into parens.
* config/i386/avx512fintrin.h (_mm512_cvt_roundps_ph, _mm512_cvtps_ph,
_mm512_mask_cvt_roundps_ph, _mm512_mask_cvtps_ph,
_mm512_maskz_cvt_roundps_ph, _mm512_maskz_cvtps_ph,
_mm512_mask_cmp_epi64_mask, _mm512_mask_cmp_epi32_mask,
_mm512_mask_cmp_epu64_mask, _mm512_mask_cmp_epu32_mask,
_mm512_mask_cmp_round_pd_mask, _mm512_mask_cmp_round_ps_mask,
_mm512_mask_cmp_pd_mask, _mm512_mask_cmp_ps_mask): Likewise.
* config/i386/avx512vlbwintrin.h (_mm256_mask_alignr_epi8,
_mm256_maskz_alignr_epi8, _mm_mask_alignr_epi8, _mm_maskz_alignr_epi8,
_mm256_mask_cmp_epu8_mask): Likewise.
* config/i386/avx512vlintrin.h (_mm_mask_cvtps_ph, _mm_maskz_cvtps_ph,
_mm256_mask_cvtps_ph, _mm256_maskz_cvtps_ph): Likewise.
* config/i386/f16cintrin.h (_mm_cvtps_ph, _mm256_cvtps_ph): Likewise.
* config/i386/shaintrin.h (_mm_sha1rnds4_epu32): Likewise.

5 years agox86: Fix -O0 intrinsic *gather*/*scatter* macros [PR94832]
Jakub Jelinek [Wed, 29 Apr 2020 15:30:22 +0000 (17:30 +0200)] 
x86: Fix -O0 intrinsic *gather*/*scatter* macros [PR94832]

As reported in the PR, while most intrinsic -O0 macro argument uses
are properly wrapped in ()s or used in context where having a complex
expression passed as the argument doesn't pose a problem (e.g. when
macro argument use is in between commas, or between ( and comma, or
between comma and ) etc.), especially the gather/scatter macros don't do
this and if one passes to some macro e.g. x + y as argument, the
corresponding inline function would do cast on the argument, but
the macro does (int) ARG, then it is (int) x + y rather than (int) (x + y).

The following patch fixes those issues in *gather/*scatter*; additionally,
the AVX2 macros were passing incorrect mask of e.g.
(__v2df)_mm_set1_pd((double)(long long int) -1)
which is IMHO equivalent to
(__v2df){-1.0, -1.0}
when it really wants to pass __v2df vector with all bits set.
I've used what the inline functions use for those cases.

2020-04-29  Jakub Jelinek  <jakub@redhat.com>

PR target/94832
* config/i386/avx2intrin.h (_mm_mask_i32gather_pd,
_mm256_mask_i32gather_pd, _mm_mask_i64gather_pd,
_mm256_mask_i64gather_pd, _mm_mask_i32gather_ps,
_mm256_mask_i32gather_ps, _mm_mask_i64gather_ps,
_mm256_mask_i64gather_ps, _mm_i32gather_epi64,
_mm_mask_i32gather_epi64, _mm256_i32gather_epi64,
_mm256_mask_i32gather_epi64, _mm_i64gather_epi64,
_mm_mask_i64gather_epi64, _mm256_i64gather_epi64,
_mm256_mask_i64gather_epi64, _mm_i32gather_epi32,
_mm_mask_i32gather_epi32, _mm256_i32gather_epi32,
_mm256_mask_i32gather_epi32, _mm_i64gather_epi32,
_mm_mask_i64gather_epi32, _mm256_i64gather_epi32,
_mm256_mask_i64gather_epi32): Surround macro parameter uses with
parens.
(_mm_i32gather_pd, _mm256_i32gather_pd, _mm_i64gather_pd,
_mm256_i64gather_pd, _mm_i32gather_ps, _mm256_i32gather_ps,
_mm_i64gather_ps, _mm256_i64gather_ps): Likewise.  Don't use
as mask vector containing -1.0 or -1.0f elts, but instead vector
with all bits set using _mm*_cmpeq_p? with zero operands.
* config/i386/avx512fintrin.h (_mm512_i32gather_ps,
_mm512_mask_i32gather_ps, _mm512_i32gather_pd,
_mm512_mask_i32gather_pd, _mm512_i64gather_ps,
_mm512_mask_i64gather_ps, _mm512_i64gather_pd,
_mm512_mask_i64gather_pd, _mm512_i32gather_epi32,
_mm512_mask_i32gather_epi32, _mm512_i32gather_epi64,
_mm512_mask_i32gather_epi64, _mm512_i64gather_epi32,
_mm512_mask_i64gather_epi32, _mm512_i64gather_epi64,
_mm512_mask_i64gather_epi64, _mm512_i32scatter_ps,
_mm512_mask_i32scatter_ps, _mm512_i32scatter_pd,
_mm512_mask_i32scatter_pd, _mm512_i64scatter_ps,
_mm512_mask_i64scatter_ps, _mm512_i64scatter_pd,
_mm512_mask_i64scatter_pd, _mm512_i32scatter_epi32,
_mm512_mask_i32scatter_epi32, _mm512_i32scatter_epi64,
_mm512_mask_i32scatter_epi64, _mm512_i64scatter_epi32,
_mm512_mask_i64scatter_epi32, _mm512_i64scatter_epi64,
_mm512_mask_i64scatter_epi64): Surround macro parameter uses with
parens.
* config/i386/avx512pfintrin.h (_mm512_prefetch_i32gather_pd,
_mm512_prefetch_i32gather_ps, _mm512_mask_prefetch_i32gather_pd,
_mm512_mask_prefetch_i32gather_ps, _mm512_prefetch_i64gather_pd,
_mm512_prefetch_i64gather_ps, _mm512_mask_prefetch_i64gather_pd,
_mm512_mask_prefetch_i64gather_ps, _mm512_prefetch_i32scatter_pd,
_mm512_prefetch_i32scatter_ps, _mm512_mask_prefetch_i32scatter_pd,
_mm512_mask_prefetch_i32scatter_ps, _mm512_prefetch_i64scatter_pd,
_mm512_prefetch_i64scatter_ps, _mm512_mask_prefetch_i64scatter_pd,
_mm512_mask_prefetch_i64scatter_ps): Likewise.
* config/i386/avx512vlintrin.h (_mm256_mmask_i32gather_ps,
_mm_mmask_i32gather_ps, _mm256_mmask_i32gather_pd,
_mm_mmask_i32gather_pd, _mm256_mmask_i64gather_ps,
_mm_mmask_i64gather_ps, _mm256_mmask_i64gather_pd,
_mm_mmask_i64gather_pd, _mm256_mmask_i32gather_epi32,
_mm_mmask_i32gather_epi32, _mm256_mmask_i32gather_epi64,
_mm_mmask_i32gather_epi64, _mm256_mmask_i64gather_epi32,
_mm_mmask_i64gather_epi32, _mm256_mmask_i64gather_epi64,
_mm_mmask_i64gather_epi64, _mm256_i32scatter_ps,
_mm256_mask_i32scatter_ps, _mm_i32scatter_ps, _mm_mask_i32scatter_ps,
_mm256_i32scatter_pd, _mm256_mask_i32scatter_pd, _mm_i32scatter_pd,
_mm_mask_i32scatter_pd, _mm256_i64scatter_ps,
_mm256_mask_i64scatter_ps, _mm_i64scatter_ps, _mm_mask_i64scatter_ps,
_mm256_i64scatter_pd, _mm256_mask_i64scatter_pd, _mm_i64scatter_pd,
_mm_mask_i64scatter_pd, _mm256_i32scatter_epi32,
_mm256_mask_i32scatter_epi32, _mm_i32scatter_epi32,
_mm_mask_i32scatter_epi32, _mm256_i32scatter_epi64,
_mm256_mask_i32scatter_epi64, _mm_i32scatter_epi64,
_mm_mask_i32scatter_epi64, _mm256_i64scatter_epi32,
_mm256_mask_i64scatter_epi32, _mm_i64scatter_epi32,
_mm_mask_i64scatter_epi32, _mm256_i64scatter_epi64,
_mm256_mask_i64scatter_epi64, _mm_i64scatter_epi64,
_mm_mask_i64scatter_epi64): Likewise.

5 years agofortran/io.c: Fix use of uninitialized variable num [PR94769]
Stefan Schulze Frielinghaus [Tue, 28 Apr 2020 11:14:28 +0000 (13:14 +0200)] 
fortran/io.c: Fix use of uninitialized variable num [PR94769]

While bootstrapping GCC on S/390 the following warning occurs:

gcc/fortran/io.c: In function 'bool gfc_resolve_dt(gfc_code*, gfc_dt*, locus*)':
gcc/fortran/io.c:3857:7: error: 'num' may be used uninitialized in this function [-Werror=maybe-uninitialized]
 3857 |       if (num == 0)
      |       ^~
gcc/fortran/io.c:3843:11: note: 'num' was declared here
 3843 |       int num;

Since gfc_resolve_dt is a non-static function we cannot assume anything about
argument DT.  Argument DT gets passed to function check_io_constraints which
passes values depending on DT, namely dt->asynchronous->value.character.string
to function compare_to_allowed_values as well as argument warn which is true as
soon as DT->dterr is true.  Thus both arguments depend on DT.

If function compare_to_allowed_values is called with
dt->asynchronous->value.character.string not being an allowed value, and
ALLOWED_F2003 as well as ALLOWED_GNU being NULL (which is the case at the
particular call side), and WARN equals true, then the function returns with a
non-zero value and leaves num uninitialized which renders the warning true.

Initialized num to -1 and added an assert statement.

gcc/fortran/ChangeLog:

2020-04-29  Stefan Schulze Frielinghaus  <stefansf@linux.ibm.com>

        PR fortran/94769
        * io.c (check_io_constraints): Initialize local variable num to
        -1 and assert that it receives a meaningful value by function
        compare_to_allowed_values.

5 years ago Fix some testsuite failures for H8/SX multilibs where short branches where used...
Jeff Law [Wed, 29 Apr 2020 14:19:22 +0000 (10:19 -0400)] 
Fix some testsuite failures for H8/SX multilibs where short branches where used when long branches were necessary.

* config/h8300/h8300.md (H8/SX div patterns): All H8/SX specific
division instructions are 4 bytes long.

5 years agors6000: Fix rs6000_atomic_assign_expand_fenv [PR94826]
Jakub Jelinek [Wed, 29 Apr 2020 13:55:39 +0000 (15:55 +0200)] 
rs6000: Fix rs6000_atomic_assign_expand_fenv [PR94826]

This is the rs6000 version of the earlier committed x86, aarch64 and arm
fixes, as create_tmp_var_raw is used because the C FE can call this outside
of function context, we need to make sure the first references to those
VAR_DECLs are through a TARGET_EXPR, so that it gets gimple_add_tmp_var
marked in whatever function it gets expanded in.  Without that DECL_CONTEXT
is NULL and the vars aren't added as local decls of the containing function.

2020-04-29  Jakub Jelinek  <jakub@redhat.com>

PR target/94826
* config/rs6000/rs6000.c (rs6000_atomic_assign_expand_fenv): Use
TARGET_EXPR instead of MODIFY_EXPR for first assignment to
fenv_var, fenv_clear and old_fenv variables.  For fenv_addr
take address of TARGET_EXPR of fenv_var with void_node initializer.
Formatting fixes.

5 years agotree-optimization: Fix use of uninitialized variable [PR94774]
Stefan Schulze Frielinghaus [Mon, 27 Apr 2020 16:09:07 +0000 (18:09 +0200)] 
tree-optimization: Fix use of uninitialized variable [PR94774]

Array retval is not necessarily initialized by function is_call_safe and
may be used afterwards.  Thus, initialize it explicitly.

gcc/ChangeLog:

2020-04-29  Stefan Schulze Frielinghaus  <stefansf@linux.ibm.com>

PR tree-optimization/94774
* gimple-ssa-sprintf.c (try_substitute_return_value): Initialize
variable retval.

5 years agoc++: Nondeterministic concepts diagnostics [PR94830]
Patrick Palka [Wed, 29 Apr 2020 13:04:58 +0000 (09:04 -0400)] 
c++: Nondeterministic concepts diagnostics [PR94830]

This patch makes the order in which template parameters appear in the
TREE_LIST returned by find_template_parameters deterministic between
runs.

The current nondeterminism is semantically harmless, but it has the
undesirable effect of causing some concepts diagnostics which print a
constraint's parameter mapping via pp_cxx_parameter_mapping to also be
nondeterministic, as in the testcases below.

gcc/cp/ChangeLog:

PR c++/94830
* pt.c (find_template_parameter_info::parm_list): New field.
(keep_template_parm): Use the new field to build up the
parameter list here instead of ...
(find_template_parameters): ... here.  Return ftpi.parm_list.

gcc/testsuite/ChangeLog:

PR c++/94830
* g++.dg/concepts/diagnostics12.C: Clarify the dg-message now
that the corresponding diagnostic is deterministic.
* g++.dg/concepts/diagnostics13.C: New test.

5 years agocalls: Remove FIXME for cxx17_empty_base_field_p
Richard Sandiford [Wed, 29 Apr 2020 12:47:21 +0000 (13:47 +0100)] 
calls: Remove FIXME for cxx17_empty_base_field_p

This predicate is now used by aarch64 targets.

2020-04-29  Richard Sandiford  <richard.sandiford@arm.com>

gcc/
* calls.h (cxx17_empty_base_field_p): Turn into a function declaration.
* calls.c (cxx17_empty_base_field_p): New function.  Check
DECL_ARTIFICIAL and RECORD_OR_UNION_TYPE_P in addition to the
previous checks.

5 years agox86: Allow -fcf-protection with external thunk
H.J. Lu [Wed, 29 Apr 2020 11:52:46 +0000 (04:52 -0700)] 
x86: Allow -fcf-protection with external thunk

Allow -fcf-protection with external thunk since the external thunk can be
made compatible with -fcf-protection.

gcc/

PR target/93654
* config/i386/i386-options.c (ix86_set_indirect_branch_type):
Allow -fcf-protection with -mindirect-branch=thunk-extern and
-mfunction-return=thunk-extern.
* doc/invoke.texi: Update notes for -fcf-protection=branch with
-mindirect-branch=thunk-extern and -mindirect-return=thunk-extern.

gcc/testsuite/

PR target/93654
* gcc.target/i386/pr93654.c: New test.

5 years agodoc: Add missing arm_arch_v8a_hard_ok anchor
Richard Sandiford [Wed, 29 Apr 2020 11:06:41 +0000 (12:06 +0100)] 
doc: Add missing arm_arch_v8a_hard_ok anchor

2020-04-29  Richard Sandiford  <richard.sandiford@arm.com>

gcc/
* doc/sourcebuild.texi: Add missing arm_arch_v8a_hard_ok anchor.