From: Thomas Schwinge <thomas@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>
Cc: Bernd Schmidt <bernds@codesourcery.com>,
Cesar Philippidis <cesar@codesourcery.com>,
Chung-Lin Tang <cltang@codesourcery.com>,
James Norris <jnorris@codesourcery.com>,
Joseph Myers <joseph@codesourcery.com>,
Julian Brown <julian@codesourcery.com>,
Tom de Vries <tom@codesourcery.com>
Subject: [gomp4] Next set of OpenACC changes
Date: Mon, 11 May 2015 16:35:00 -0000 [thread overview]
Message-ID: <87mw1bkr7j.fsf@schwinge.name> (raw)
In-Reply-To: <87sibbpfpx.fsf@schwinge.name>
[-- Attachment #1: Type: text/plain, Size: 30715 bytes --]
Hi!
On Tue, 05 May 2015 10:54:02 +0200, I wrote:
> In follow-up messages, I'll be posting the separated parts (for easier
> review) of a next set of OpenACC changes that we'd like to commit.
> ChangeLog updates not yet written; will do that before commit, obviously.
In order for us to be able to make progress with staging our other
OpenACC changes in gomp-4_0-branch, I have now committed to
gomp-4_0-branch r223007, which is these patches as posted plus a tiny
last-minute typo fix (see below), and we shall then work on addressing
the review comments already provided (thanks!) (as well as those which I
found myself, upon reviewing our changes), before later re-submitting for
trunk.
commit cd00a35cd24a3ac05cac0061639ce631e52f2f49
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon May 11 16:29:03 2015 +0000
Next set of OpenACC changes
gcc/c-family/
* c-common.c (c_common_attribute_table): Set min_len to -1 for
"omp declare target". Add "oacc declare".
* c-common.h (oacc_extract_device_id, oacc_filter_device_types):
New prototypes.
* c-omp.c (oacc_extract_device_id, oacc_filter_device_types): New
functions.
* c-pragma.c (oacc_pragmas): Add "atomic", "declare", "host_data",
"routine".
* c-pragma.h (pragma_kind): Add PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_DECLARE, PRAGMA_OACC_HOST_DATA, PRAGMA_OACC_ROUTINE.
(pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_LINK, PRAGMA_OACC_CLAUSE_NOHOST,
PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE,
PRAGMA_OACC_CLAUSE_DEFAULT, remove PRAGMA_OACC_CLAUSE_SELF.
gcc/c/
* c-parser.c (c_parser): Add oacc_routines member.
(c_parse_file): Initialize it.
(c_parser_declaration_or_fndef): Add oacc_routine_clauses, and
oacc_routine_named formal parameters. Adjust all users. Support
OpenACC routines.
(c_parser_pragma): Handle PRAGMA_OACC_DECLARE,
PRAGMA_OACC_ROUTINE, PRAGMA_OACC_WAIT. Add pragma context
checking for PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA.
(c_parser_omp_clause_name): Add consume_token formal parameter.
Handle "bind", "device_resident", "device_type", "dtype",
"independent", "link", "nohost", "tile", "use_device".
(c_parser_oacc_wait_list): Change an error message.
(c_parser_oacc_data_clause): Handle
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_LINK.
Don't handle PRAGMA_OACC_CLAUSE_SELF.
(c_parser_omp_clause_default): Add only_none formal parameter.
(c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_threads)
(c_parser_omp_clause_num_workers)
(c_parser_omp_clause_vector_length): Replace functions by...
(require_positive_expr, c_parser_omp_positive_int_clause):
... these new functions. Adjust all users.
(c_parser_omp_clause_untied, c_parser_omp_clause_branch): Replace
functions by...
(c_parser_omp_simple_clause): ... this new function. Adjust all
users.
(c_parser_oacc_shape_clause, c_parser_oacc_clause_bind)
(c_parser_oacc_clause_device_type, c_parser_oacc_clause_tile)
(c_parser_oacc_clause_use_device): New functions.
(c_parser_oacc_all_clauses): Add dtype_mask, and scan_dtype formal
parameters. Adjust all users. Handle PRAGMA_OACC_CLAUSE_AUTO,
PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OMP_CLAUSE_DEFAULT,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_LINK,
PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_SEQ,
PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE,
PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_WORKER. Don't
handle PRAGMA_OACC_CLAUSE_SELF.
(c_parser_oacc_declare, oacc_split_loop_clauses)
(c_parser_oacc_host_data, c_parser_oacc_routine)
(c_finish_oacc_routine): New functions.
(c_parser_oacc_enter_exit_data): Change error reporting.
(c_parser_oacc_loop): Add mask, and cclauses formal parameters.
Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_WORKER, PRAGMA_OACC_CLAUSE_VECTOR,
PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_PRIVATE.
(c_parser_oacc_kernels): Support combined directives. Handle
PRAGMA_OACC_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE.
(c_parser_oacc_parallel): Likewise. Handle
PRAGMA_OACC_CLAUSE_PRIVATE.
(c_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_HOST_DATA.
(c_parser_oacc_update): Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE,
PRAGMA_OACC_CLAUSE_WAIT. Don't handle PRAGMA_OACC_CLAUSE_SELF.
* c-tree.h (c_finish_oacc_host_data): New prototype.
* c-typeck.c (c_finish_oacc_host_data): New function.
(c_finish_omp_clauses): Add oacc formal parameter. Adjust all
users. Handle OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_USE_DEVICE,
OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE.
gcc/cp/
* cp-gimplify.c (cxx_omp_clause_copy_ctor): Handle OMP_CLAUSE_MAP.
* cp-tree.h (finish_oacc_host_data): New prototype.
* parser.h (cp_parser): Add oacc_routine, named_oacc_routines
members.
* parser.c (cp_parser_new): Initialize them.
(cp_ensure_no_omp_declare_simd, cp_parser_init_declarator)
(cp_parser_late_return_type_opt, cp_parser_member_declaration)
(cp_parser_function_definition_from_specifiers_and_declarator)
(cp_parser_save_member_function_body, cp_parser_omp_declare_simd)
(cp_parser_omp_declare, cp_parser_pragma): Extend for OpenACC routines.
(cp_finalize_oacc_routine)
(cp_parser_oacc_routine_check_parallelism, cp_parser_oacc_routine)
(cp_parser_late_parsing_oacc_routine): New functions.
(cp_parser_omp_clause_name): Add consume_token formal parameter.
Handle "auto", "bind", "device_resident", "device_type", "dtype",
"gang", "independent", "link", "nohost", "seq", "tile",
"use_device", "vector", "worker".
(cp_parser_oacc_data_clause): Handle
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_LINK.
Don't handle PRAGMA_OACC_CLAUSE_SELF.
(cp_parser_oacc_clause_vector_length)
(cp_parser_omp_clause_num_gangs, cp_parser_omp_clause_num_threads)
(cp_parser_omp_clause_num_workers): Replace functions by...
(require_positive_expr, cp_parser_omp_positive_int_clause):
... these new functions. Adjust all users.
(cp_parser_oacc_shape_clause, cp_parser_oacc_clause_device_type)
(cp_parser_oacc_clause_tile, cp_parser_oacc_clause_bind): New
functions.
(cp_parser_oacc_wait_list): Change an error message.
(cp_parser_omp_clause_default): Add is_omp formal parameter.
(cp_parser_omp_clause_untied, cp_parser_omp_clause_branch):
Replace functions by...
(cp_parser_omp_simple_clause): ... this new function. Adjust all
users.
(cp_parser_oacc_all_clauses): Add dtype_mask, and scan_dtype
formal parameters. Adjust all users. Handle
PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_BIND,
PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_LINK,
PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_PRIVATE,
PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_VECTOR,
PRAGMA_OACC_CLAUSE_WORKER. Don't handle PRAGMA_OACC_CLAUSE_SELF.
(cp_parser_oacc_host_data, oacc_split_loop_clauses):
(cp_parser_oacc_enter_exit_data): Change error reporting.
(cp_parser_oacc_loop): Add p_name, mask, and cclauses formal
parameters. Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE,
PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_PRIVATE
PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_WORKER,
PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE. Support combined
directives.
(cp_parser_oacc_kernels, cp_parser_oacc_parallel): Replace
functions by...
(cp_parser_oacc_parallel_kernels): ... this new function. Adjust
all users. Support combined directives. For "kernels", handle
PRAGMA_OACC_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_DEVICE_TYPE. For
"parallel", handle PRAGMA_OACC_CLAUSE_DEFAULT,
PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_PRIVATE.
(cp_parser_oacc_update): Handle PRAGMA_OACC_CLAUSE_DEVICE_TYPE.
Don't handle PRAGMA_OACC_CLAUSE_SELF.
(cp_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_HOST_DATA.
(cp_parser_pragma): Handle PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_HOST_DATA, PRAGMA_OACC_ROUTINE. Add pragma context
checking for PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NUM_GANGS,
OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_GANG,
OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_ASYNC,
OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
OMP_CLAUSE_SEQ, OMP_CLAUSE_TILE.
(tsubst_expr): Handle OACC_PARALLEL, OACC_KERNELS, OACC_LOOP,
OACC_DATA, OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE.
* semantics.c (finish_omp_clauses): Add oacc formal parameter.
Adjust all users. Handle OMP_CLAUSE_GANG, OMP_CLAUSE_VECTOR,
OMP_CLAUSE_WORKER, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_AUTO, OMP_CLAUSE_INDEPENDENT,
OMP_CLAUSE_SEQ, OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST,
OMP_CLAUSE_TILE.
(finish_oacc_host_data): New function.
gcc/fortran/
* dump-parse-tree.c (show_namespace): Rewrite handling of OpenACC
declare.
* gfortran.h (gfc_statement): Add ST_OACC_ATOMIC,
ST_OACC_END_ATOMIC.
(gfc_omp_clauses): Add routine_bind, dtype, dtype_clauses, nohost,
acc_collapse, bind, num_gangs, num_workers, vector_length, tile
members.
(gfc_oacc_declare, gfc_oacc_routine_name): New typedefs.
(gfc_get_oacc_declare, gfc_get_oacc_routine_name): New macros.
(gfc_namespace): Add oacc_declare, oacc_routine_clauses,
oacc_routine_names, oacc_routine members, remove
oacc_declare_clauses member.
(gfc_exec_op): Add EXEC_OACC_ROUTINE, EXEC_OACC_ATOMIC,
EXEC_OACC_DECLARE.
(gfc_code): Add oacc_declare member.
(gfc_free_oacc_declares, insert_oacc_declare): New prototypes.
* match.h (gfc_match_oacc_atomic): New prototype.
* openmp.c (OMP_CLAUSE_HOST_SELF): Rename to...
(OMP_CLAUSE_HOST): ... this. Adjust all users.
(OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST, OMP_CLAUSE_DEVICE_TYPE): New
macros.
(gfc_free_oacc_declares): New function.
(gfc_match_omp_map_clause): Add allow_sections formal parameter.
Adjust all users.
(gfc_match_omp_clauses): Add dtype_mask formal parameter. Adjust
all users. Change handling of OMP_CLAUSE_VECTOR_LENGTH,
OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_TILE,
OMP_CLAUSE_DEFAULT, OMP_CLAUSE_COLLAPSE. Handle OMP_CLAUSE_BIND,
OMP_CLAUSE_NOHOST, OMP_CLAUSE_DEVICE_TYPE.
(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_LOOP_CLAUSES)
(OACC_UPDATE_CLAUSES): Add OMP_CLAUSE_DEVICE_TYPE.
(OACC_ROUTINE_CLAUSES, OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK)
(OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK)
(OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK)
(OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK)
(OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK): New macros.
(gfc_match_oacc_declare, gfc_match_oacc_routine): Rewrite
functions.
(gfc_match_oacc_update): Add error reporting.
(gfc_match_omp_atomic, gfc_match_oacc_atomic): New wrapper
functions around...
(gfc_match_omp_oacc_atomic): ... this new function.
(check_array_not_assumed): Remove pointer check.
(oacc_code_to_statement): Handle EXEC_OACC_ATOMIC.
(resolve_oacc_loop_blocks): Don't error out for combined OpenACC
gang, worker, and vector clauses.
(resolve_oacc_cache): Remove function.
(gfc_resolve_oacc_declare): Rewrite function.
(gfc_resolve_oacc_directive): Handle EXEC_OACC_ATOMIC. Don't
handle EXEC_OACC_CACHE.
* parse.c (decode_oacc_directive): Handle "atomic", "end atomic".
(case_exec_markers): Add ST_OACC_ATOMIC.
(case_decl): Add ST_OACC_DECLARE.
(gfc_ascii_statement): Handle ST_OACC_ATOMIC, ST_OACC_END_ATOMIC.
(verify_st_order, parse_spec): Remove handling of ST_OACC_DECLARE.
(parse_omp_atomic): Rename to...
(parse_omp_oacc_atomic): ... this new function. Add omp_p formal
parameter. Adjust all users.
(parse_executable): Handle ST_OACC_ATOMIC.
(parse_progunit): Remove handling of OpenACC declare.
(is_oacc): Handle EXEC_OACC_ROUTINE.
* parse.h (gfc_state_data): Add ext.oacc_declare member. Remove
ext.oacc_declare_clauses member.
* resolve.c (gfc_resolve_blocks): Handle EXEC_OACC_ATOMIC,
EXEC_OACC_ROUTINE, EXEC_OACC_DECLARE.
(gfc_resolve_code): Handle EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE.
* st.c (gfc_free_statement): Handle EXEC_OACC_DECLARE,
EXEC_OACC_ROUTINE, EXEC_OACC_ATOMIC.
* trans-decl.c (find_end, insert_oacc_declare): New functions.
(gfc_generate_function_code): Change handling of OpenACC declare.
* trans-openmp.c (gfc_omp_clause_copy_ctor): Handle
OMP_CLAUSE_REDUCTION.
(gfc_trans_omp_clauses): Add appropriate, generate OMP_CLAUSE_SEQ
(instead of OMP_CLAUSE_ORDERED), OMP_CLAUSE_AUTO, or
OMP_CLAUSE_TILE.
(gfc_trans_oacc_combined_directive): Don't set
OACC_KERNELS_COMBINED, and OACC_PARALLEL_COMBINED.
(gfc_trans_oacc_declare): Rewrite function.
(gfc_trans_oacc_directive): Handle EXEC_OACC_ATOMIC,
EXEC_OACC_DECLARE.
* trans-stmt.c (gfc_trans_block_construct): Change handling of
OpenACC declare.
* trans.c (trans_code): Handle EXEC_OACC_ATOMIC,
EXEC_OACC_DECLARE.
gcc/
* gimplify.c (gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_TILE.
(gimplify_expr): Don't verify OACC_KERNELS_COMBINED, and
OACC_PARALLEL_COMBINED.
* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_BIND,
OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE.
(check_omp_nesting_restrictions): Support GIMPLE_OMP_ATOMIC_LOAD,
GIMPLE_OMP_ATOMIC_STORE inside OpenACC contexts.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE_BIND,
OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE, OMP_CLAUSE_DEVICE_TYPE.
* 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_TILE.
* tree.h (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): Remove
macros.
(OMP_CLAUSE_BIND_NAME, OMP_CLAUSE_TILE_LIST)
(OMP_CLAUSE_DEVICE_TYPE_DEVICES, OMP_CLAUSE_DEVICE_TYPE_CLAUSES):
Add macros.
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-1.c: Update.
* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
* c-c++-common/goacc/asyncwait-1.c: Likewise.
* c-c++-common/goacc/data-2.c: Likewise.
* c-c++-common/goacc/reduction-1.c: Likewise.
* c-c++-common/goacc/reduction-2.c: Likewise.
* c-c++-common/goacc/reduction-3.c: Likewise.
* c-c++-common/goacc/reduction-4.c: Likewise.
* gfortran.dg/goacc/cache-1.f95: Likewise.
* gfortran.dg/goacc/coarray.f95: Likewise.
* gfortran.dg/goacc/coarray_2.f90: Likewise.
* gfortran.dg/goacc/combined_loop.f90: Likewise.
* gfortran.dg/goacc/cray.f95: Likewise.
* gfortran.dg/goacc/declare-1.f95: Likewise.
* gfortran.dg/goacc/host_data-tree.f95: Likewise.
* gfortran.dg/goacc/loop-1.f95: Likewise.
* gfortran.dg/goacc/loop-2.f95: Likewise.
* gfortran.dg/goacc/parameter.f95: Likewise.
* c-c++-common/goacc/loop-1.c: Enable for C++.
* c-c++-common/goacc/kernels-1.c: Rename to...
* c-c++-common/goacc/kernels-empty.c: ... this new file.
* c-c++-common/goacc/parallel-1.c: Rename to...
* c-c++-common/goacc/parallel-empty.c: ... this new file.
* c-c++-common/goacc/declare-1.c: New file.
* c-c++-common/goacc/declare-2.c: Likewise.
* c-c++-common/goacc/dtype-1.c: Likewise.
* c-c++-common/goacc/dtype-2.c: Likewise.
* c-c++-common/goacc/host_data-1.c: Likewise.
* c-c++-common/goacc/host_data-2.c: Likewise.
* c-c++-common/goacc/host_data-3.c: Likewise.
* c-c++-common/goacc/host_data-4.c: Likewise.
* c-c++-common/goacc/kernels-eternal.c: Likewise.
* c-c++-common/goacc/kernels-noreturn.c: Likewise.
* c-c++-common/goacc/parallel-eternal.c: Likewise.
* c-c++-common/goacc/parallel-noreturn.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-2.c: Likewise.
* c-c++-common/goacc/routine-3.c: Likewise.
* c-c++-common/goacc/routine-4.c: Likewise.
* c-c++-common/goacc/tile.c: Likewise.
* g++.dg/goacc/template-reduction.C: Likewise.
* g++.dg/goacc/template.C: Likewise.
* gfortran.dg/goacc/declare-2.f95: Likewise.
* gfortran.dg/goacc/default.f95: Likewise.
* gfortran.dg/goacc/dtype-1.f95: Likewise.
* gfortran.dg/goacc/dtype-2.f95: Likewise.
* gfortran.dg/goacc/modules.f95: Likewise.
* gfortran.dg/goacc/update.f95: Likewise.
include/
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT,
GOMP_MAP_LINK.
libgomp/
* oacc-mem.c (update_dev_host): Add missing initialization.
* oacc-ptx.h (GOMP_ATOMIC_PTX): New macro.
* plugin/plugin-nvptx.c (link_ptx): Link it in.
* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Update.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-71.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-72.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-74.c: Likewise.
* 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-83.c: Likewise.
* testsuite/libgomp.oacc-fortran/data-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-3.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-4.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/subr.h: Likewise.
* testsuite/libgomp.oacc-c-c++-common/subr.ptx: Likewise.
* testsuite/libgomp.oacc-c-c++-common/timer.h: Remove file.
* testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Move common
code from here...
* testsuite/libgomp.oacc-c-c++-common/parallel-1.c: ..., and
here...
* testsuite/libgomp.oacc-c-c++-common/data-clauses.h: ... into
this new file.
* testsuite/libgomp.oacc-c++/template-reduction.C: New test.
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-2.c: Likewise.
* testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/atomic_update-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/cache-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/clauses-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/declare-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise.
* testsuite/libgomp.oacc-fortran/lib-13.f90: Likewise.
* testsuite/libgomp.oacc-fortran/lib-14.f90: Likewise.
* testsuite/libgomp.oacc-fortran/lib-15.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-5.f90: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@223007 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 28 +
gcc/c-family/ChangeLog.gomp | 26 +
gcc/c-family/c-common.c | 3 +-
gcc/c-family/c-common.h | 2 +
gcc/c-family/c-omp.c | 105 ++
gcc/c-family/c-pragma.c | 4 +
gcc/c-family/c-pragma.h | 14 +-
gcc/c/ChangeLog.gomp | 71 +
gcc/c/c-parser.c | 1353 ++++++++++++----
gcc/c/c-tree.h | 3 +-
gcc/c/c-typeck.c | 112 +-
gcc/cp/ChangeLog.gomp | 93 ++
gcc/cp/cp-gimplify.c | 3 +-
gcc/cp/cp-tree.h | 3 +-
gcc/cp/parser.c | 1382 +++++++++++++----
gcc/cp/parser.h | 4 +
gcc/cp/pt.c | 43 +-
gcc/cp/semantics.c | 151 +-
gcc/fortran/ChangeLog.gomp | 94 ++
gcc/fortran/dump-parse-tree.c | 12 +-
gcc/fortran/gfortran.h | 50 +-
gcc/fortran/match.h | 1 +
gcc/fortran/openmp.c | 581 +++++--
gcc/fortran/parse.c | 65 +-
gcc/fortran/parse.h | 2 +-
gcc/fortran/resolve.c | 5 +
gcc/fortran/st.c | 7 +
gcc/fortran/trans-decl.c | 62 +-
gcc/fortran/trans-openmp.c | 66 +-
gcc/fortran/trans-stmt.c | 7 +-
gcc/fortran/trans-stmt.h | 2 +-
gcc/fortran/trans.c | 2 +
gcc/gimplify.c | 16 +-
gcc/omp-low.c | 11 +-
gcc/testsuite/ChangeLog.gomp | 58 +
gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c | 46 +
.../c-c++-common/goacc-gomp/nesting-fail-1.c | 25 -
gcc/testsuite/c-c++-common/goacc/asyncwait-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/data-2.c | 12 +-
gcc/testsuite/c-c++-common/goacc/declare-1.c | 84 +
gcc/testsuite/c-c++-common/goacc/declare-2.c | 67 +
gcc/testsuite/c-c++-common/goacc/dtype-1.c | 113 ++
gcc/testsuite/c-c++-common/goacc/dtype-2.c | 31 +
gcc/testsuite/c-c++-common/goacc/host_data-1.c | 14 +
gcc/testsuite/c-c++-common/goacc/host_data-2.c | 14 +
gcc/testsuite/c-c++-common/goacc/host_data-3.c | 16 +
gcc/testsuite/c-c++-common/goacc/host_data-4.c | 15 +
.../goacc/{kernels-1.c => kernels-empty.c} | 0
gcc/testsuite/c-c++-common/goacc/kernels-eternal.c | 11 +
.../c-c++-common/goacc/kernels-noreturn.c | 12 +
gcc/testsuite/c-c++-common/goacc/loop-1.c | 2 -
.../goacc/{parallel-1.c => parallel-empty.c} | 0
.../c-c++-common/goacc/parallel-eternal.c | 11 +
.../c-c++-common/goacc/parallel-noreturn.c | 12 +
gcc/testsuite/c-c++-common/goacc/reduction-1.c | 25 +-
gcc/testsuite/c-c++-common/goacc/reduction-2.c | 22 +-
gcc/testsuite/c-c++-common/goacc/reduction-3.c | 22 +-
gcc/testsuite/c-c++-common/goacc/reduction-4.c | 40 +-
gcc/testsuite/c-c++-common/goacc/routine-1.c | 35 +
gcc/testsuite/c-c++-common/goacc/routine-2.c | 36 +
gcc/testsuite/c-c++-common/goacc/routine-3.c | 52 +
gcc/testsuite/c-c++-common/goacc/routine-4.c | 87 ++
gcc/testsuite/c-c++-common/goacc/tile.c | 26 +
gcc/testsuite/g++.dg/goacc/template-reduction.C | 100 ++
gcc/testsuite/g++.dg/goacc/template.C | 131 ++
gcc/testsuite/gfortran.dg/goacc/cache-1.f95 | 1 -
gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 2 +-
gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 | 1 +
gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 | 2 +-
gcc/testsuite/gfortran.dg/goacc/cray.f95 | 1 -
gcc/testsuite/gfortran.dg/goacc/declare-1.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 44 +
gcc/testsuite/gfortran.dg/goacc/default.f95 | 17 +
gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 161 ++
gcc/testsuite/gfortran.dg/goacc/dtype-2.f95 | 39 +
gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 2 +-
gcc/testsuite/gfortran.dg/goacc/loop-1.f95 | 1 -
gcc/testsuite/gfortran.dg/goacc/loop-2.f95 | 26 +-
gcc/testsuite/gfortran.dg/goacc/modules.f95 | 55 +
gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 1 -
gcc/testsuite/gfortran.dg/goacc/update.f95 | 5 +
gcc/tree-core.h | 14 +-
gcc/tree-pretty-print.c | 6 +
gcc/tree.c | 13 +-
gcc/tree.h | 21 +-
include/ChangeLog.gomp | 6 +
include/gomp-constants.h | 4 +
libgomp/ChangeLog.gomp | 70 +
libgomp/oacc-mem.c | 3 +
libgomp/oacc-ptx.h | 28 +
libgomp/plugin/plugin-nvptx.c | 10 +
.../libgomp.oacc-c++/template-reduction.C | 102 ++
.../libgomp.oacc-c-c++-common/atomic_capture-1.c | 866 +++++++++++
.../libgomp.oacc-c-c++-common/atomic_capture-2.c | 1626 ++++++++++++++++++++
.../libgomp.oacc-c-c++-common/atomic_update-1.c | 760 +++++++++
.../libgomp.oacc-c-c++-common/clauses-1.c | 26 +
.../testsuite/libgomp.oacc-c-c++-common/data-2.c | 44 +-
.../testsuite/libgomp.oacc-c-c++-common/data-3.c | 18 +-
.../{parallel-1.c => data-clauses.h} | 32 +-
.../libgomp.oacc-c-c++-common/kernels-1.c | 182 +--
.../testsuite/libgomp.oacc-c-c++-common/lib-69.c | 70 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-70.c | 79 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-71.c | 55 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-72.c | 60 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-73.c | 64 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-74.c | 91 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-75.c | 89 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-76.c | 88 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-77.c | 91 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-78.c | 91 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-79.c | 91 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-80.c | 95 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-81.c | 106 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-82.c | 43 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-83.c | 22 +-
.../libgomp.oacc-c-c++-common/parallel-1.c | 204 +--
.../libgomp.oacc-c-c++-common/routine-1.c | 40 +
.../libgomp.oacc-c-c++-common/routine-2.c | 41 +
libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h | 44 +-
.../testsuite/libgomp.oacc-c-c++-common/subr.ptx | 222 +--
.../testsuite/libgomp.oacc-c-c++-common/timer.h | 103 --
.../libgomp.oacc-fortran/atomic_capture-1.f90 | 784 ++++++++++
.../libgomp.oacc-fortran/atomic_update-1.f90 | 338 ++++
libgomp/testsuite/libgomp.oacc-fortran/cache-1.f90 | 26 +
.../testsuite/libgomp.oacc-fortran/clauses-1.f90 | 290 ++++
libgomp/testsuite/libgomp.oacc-fortran/data-1.f90 | 231 ++-
libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 | 50 +
libgomp/testsuite/libgomp.oacc-fortran/data-3.f90 | 34 +-
.../testsuite/libgomp.oacc-fortran/data-4-2.f90 | 19 +-
libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 | 19 +-
.../testsuite/libgomp.oacc-fortran/declare-1.f90 | 229 +++
libgomp/testsuite/libgomp.oacc-fortran/lib-12.f90 | 24 +
libgomp/testsuite/libgomp.oacc-fortran/lib-13.f90 | 28 +
libgomp/testsuite/libgomp.oacc-fortran/lib-14.f90 | 79 +
libgomp/testsuite/libgomp.oacc-fortran/lib-15.f90 | 52 +
.../testsuite/libgomp.oacc-fortran/routine-5.f90 | 27 +
136 files changed, 11216 insertions(+), 2501 deletions(-)
This is the patches as posted, including the following last-minute typo
fix:
--- a/gcc/testsuite/c-c++-common/goacc/host_data-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-4.c
@@ -6,7 +6,7 @@ int main (int argc, char* argv[])
#pragma acc enter data copyin (x)
/* Specifying an array index is not valid for host_data/use_device. */
- #pragma acc host_data use_device (x[4]) /* { dg-error "expected \\\')' before '\\\[' token" } */
+ #pragma acc host_data use_device (x[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
;
#pragma acc exit data delete (x)
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
next prev parent reply other threads:[~2015-05-11 16:35 UTC|newest]
Thread overview: 11+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-05-05 8:54 Thomas Schwinge
2015-05-05 8:56 ` Next set of OpenACC changes: middle end, libgomp Thomas Schwinge
2015-05-05 8:58 ` Next set of OpenACC changes: C family Thomas Schwinge
2015-05-05 14:19 ` Jakub Jelinek
2015-05-05 15:40 ` Cesar Philippidis
2015-05-05 8:59 ` Next set of OpenACC changes: Fortran Thomas Schwinge
2015-05-05 10:42 ` Bernhard Reutner-Fischer
2015-05-05 9:00 ` Next set of OpenACC changes: Testsuite Thomas Schwinge
2015-05-11 16:35 ` Thomas Schwinge [this message]
2015-05-13 20:57 ` [gomp4] Assorted OpenACC changes (was: Next set of OpenACC changes) Thomas Schwinge
2015-05-14 8:37 ` Jakub Jelinek
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=87mw1bkr7j.fsf@schwinge.name \
--to=thomas@codesourcery.com \
--cc=bernds@codesourcery.com \
--cc=cesar@codesourcery.com \
--cc=cltang@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=jnorris@codesourcery.com \
--cc=joseph@codesourcery.com \
--cc=julian@codesourcery.com \
--cc=tom@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).