Commit Graph

244 Commits

Author SHA1 Message Date
Tom de Vries a6860be9c8 Add gcc.dg/parloops-exit-first-loop-alt-4.c
2015-06-13  Tom de Vries  <tom@codesourcery.com>

	* gcc.dg/parloops-exit-first-loop-alt-4.c: New test.

	* testsuite/libgomp.c/parloops-exit-first-loop-alt-4.c: New test.

From-SVN: r224453
2015-06-13 08:05:54 +00:00
Tom de Vries ab3306d23c Update parloops-exit-first-loop-alt{,-2,-3}.c
2015-06-13  Tom de Vries  <tom@codesourcery.com>

	* testsuite/libgomp.c/parloops-exit-first-loop-alt-2.c: Add comment.
	* testsuite/libgomp.c/parloops-exit-first-loop-alt.c: Same.
	* testsuite/libgomp.c/parloops-exit-first-loop-alt-3.c: Add comment.
	(N): Define.
	(main): Use N instead of hardcoded constants.

	* gcc.dg/parloops-exit-first-loop-alt-2.c: Add comment.
	(main): Remove superfluous attributes.
	* gcc.dg/parloops-exit-first-loop-alt-3.c: Same.
	* gcc.dg/parloops-exit-first-loop-alt.c: Same.

From-SVN: r224452
2015-06-13 08:05:44 +00:00
Tom de Vries 7c82d82716 Add transform_to_exit_first_loop_alt
2015-06-05  Tom de Vries  <tom@codesourcery.com>

	merge from gomp4 branch:
	2015-05-28  Tom de Vries  <tom@codesourcery.com>

	PR tree-optimization/65443
	* tree-parloops.c (replace_imm_uses, replace_uses_in_bb_by)
	(replace_uses_in_bbs_by, transform_to_exit_first_loop_alt)
	(try_transform_to_exit_first_loop_alt): New function.
	(transform_to_exit_first_loop): Use
	try_transform_to_exit_first_loop_alt.

	* gcc.dg/parloops-exit-first-loop-alt-2.c: New test.
	* gcc.dg/parloops-exit-first-loop-alt-3.c: New test.
	* gcc.dg/parloops-exit-first-loop-alt.c: New test.

	* testsuite/libgomp.c/parloops-exit-first-loop-alt-2.c: New test.
	* testsuite/libgomp.c/parloops-exit-first-loop-alt-3.c: New test.
	* testsuite/libgomp.c/parloops-exit-first-loop-alt.c: New test.

From-SVN: r224154
2015-06-05 15:57:34 +00:00
Bernhard Reutner-Fischer c469078de7 testsuite: auto-wipe dump files
gcc/testsuite/ChangeLog

2015-05-29  Bernhard Reutner-Fischer  <aldot@gcc.gnu.org>

	* lib/gcc-dg.exp (cleanup-ipa-dump, cleanup-rtl-dump,
	cleanup-tree-dump, cleanup-dump, cleanup-saved-temps): Remove.
	Adjust all callers.
	(schedule-cleanups, dg-keep-saved-temps): New proc.
	(gcc-dg-test-1): Schedule cleanups.
	* lib/profopt.exp (profopt-execute): Likewise.
	* g++.dg/cdce3.C: Adjust expected line numbers.
	* gcc.dg/cdce1.c: Likewise.
	* gcc.dg/cdce2.c: Likewise.
	* gcc.dg/strlenopt-22.c: Fix comment delimiter.
	* gcc.dg/strlenopt-24.c: Likewise.
	* gcc.dg/tree-ssa/vrp26.c: Likewise.
	* gcc.dg/tree-ssa/vrp28.c: Likewise.
	* obj-c++.dg/encode-2.mm: Likewise.

libgomp/ChangeLog

2015-05-29  Bernhard Reutner-Fischer  <aldot@gcc.gnu.org>

	* testsuite/libgomp.graphite/bounds.c: Adjust for
	cleanup-tree-dump removal.
	* testsuite/libgomp.graphite/force-parallel-1.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-2.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-3.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-4.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-5.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-6.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-7.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-8.c: Likewise.
	* testsuite/libgomp.graphite/force-parallel-9.c: Likewise.
	* testsuite/libgomp.graphite/pr41118.c: Likewise.


gcc/ChangeLog

2015-05-29  Bernhard Reutner-Fischer  <aldot@gcc.gnu.org>

	* config/arm/neon-testgen.ml (emit_epilogue): Remove manual call
	to cleanup-saved-temps.
	* doc/sourcebuild.texi (Clean up generated test files): Expand
	introduction.
	(dg-keep-saved-temps): Document new proc.
	(cleanup-ipa-dump, cleanup-rtl-dump, cleanup-tree-dump,
	cleanup-saved-temps): Remove.

From-SVN: r223858
2015-05-29 10:20:29 +02:00
Julian Brown 9fb5fd444b oacc-init.c (resolve_device): Add FAIL_IS_ERROR argument.
* oacc-init.c (resolve_device): Add FAIL_IS_ERROR argument. Update
	function comment. Only call gomp_fatal if new argument is true.
	(acc_dev_num_out_of_range): New function.
	(acc_init_1, acc_shutdown_1): Update call to resolve_device. Call
	acc_dev_num_out_of_range as appropriate.
	(acc_get_num_devices, acc_set_device_type, acc_get_device_type)
	(acc_get_device_num, acc_set_device_num): Update calls to resolve_device.
	* testsuite/libgomp.oacc-c-c++-common/lib-4.c: Update expected test
	output.

From-SVN: r223803
2015-05-28 10:35:39 +00:00
Jakub Jelinek 41b37d5ec1 re PR middle-end/66199 (lastprivate/linear clause issues on combined constructs)
PR middle-end/66199
	* tree.h (OMP_TEAMS_COMBINED): Define.
	* gimplify.c (enum gimplify_omp_var_data): Add
	GOVD_LINEAR_LASTPRIVATE_NO_OUTER.
	(enum omp_region_type): Add ORT_COMBINED_TEAMS.
	(omp_notice_variable): Accept both ORT_TEAMS
	and ORT_COMBINED_TEAMS.  Don't recurse if
	GOVD_LINEAR_LASTPRIVATE_NO_OUTER is set and either
	GOVD_LINEAR is set, or GOVD_LASTPRIVATE without
	GOVD_FIRSTPRIVATE.
	(omp_no_lastprivate): New function.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_LASTPRIVATE
	and OMP_CLAUSE_LINEAR, if omp_no_lastprivate, don't
	notice_outer and set appropriate bits, otherwise make
	sure default(none) combined constructs won't complain.
	(gimplify_adjust_omp_clauses): Remove OMP_CLAUSE_LINEAR
	outer special casing, for OMP_CLAUSE_LASTPRIVATE if
	omp_no_lastprivate either remove the clause or turn it
	into OMP_CLAUSE_PRIVATE.
	(gimplify_omp_for): Fix up handling of implicit
	lastprivate or linear iterators.
	(gimplify_omp_workshare): For OMP_TEAMS_COMBINED use
	ORT_COMBINED_TEAMS.
	* omp-low.c (lower_omp_for_lastprivate): For combined
	for simd use fd.loop.n2 from the for rather than simd.
gcc/c/
	* c-parser.c (c_parser_omp_for_loop): Don't add
	OMP_CLAUSE_SHARED to OMP_PARALLEL_CLAUSES when moving
	OMP_CLAUSE_LASTPRIVATE clause to OMP_FOR_CLAUSES.
	(c_parser_omp_teams): Set OMP_TEAMS_COMBINED for combined
	constructs.
gcc/cp/
	* parser.c (cp_parser_omp_for_loop): Don't add
	OMP_CLAUSE_SHARED to OMP_PARALLEL_CLAUSES when moving
	OMP_CLAUSE_LASTPRIVATE clause to OMP_FOR_CLAUSES.
	(cp_parser_omp_teams): Set OMP_TEAMS_COMBINED for combined
	constructs.
gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_teams): Set OMP_TEAMS_COMBINED for
	combined constructs.
	(gfc_trans_omp_target): Make sure BIND_EXPR has non-NULL
	BIND_EXPR_BLOCK.
libgomp/
	* testsuite/libgomp.c/pr66199-1.c: New test.
	* testsuite/libgomp.c/pr66199-2.c: New test.
	* testsuite/libgomp.c++/pr66199-1.C: New test.
	* testsuite/libgomp.c++/pr66199-2.C: New test.
	* testsuite/libgomp.fortran/pr66199-1.f90: New test.
	* testsuite/libgomp.fortran/pr66199-2.f90: New test.

From-SVN: r223387
2015-05-19 18:16:15 +02:00
Jakub Jelinek b37dddbc77 re PR middle-end/66133 (Wrong-code with noreturn #pragma omp task body)
PR middle-end/66133
	* omp-low.c (expand_omp_taskreg): For GIMPLE_OMP_TASK expansion,
	make sure it is never noreturn, even when the task body does not
	return.
	(lower_omp_taskreg): For GIMPLE_OMP_TASK, emit GIMPLE_OMP_CONTINUE
	right before GIMPLE_OMP_RETURN.
	(make_gimple_omp_edges): Accept GIMPLE_OMP_CONTINUE as ->cont
	for GIMPLE_OMP_TASK.  For GIMPLE_OMP_RETURN corresponding to
	GIMPLE_OMP_TASK add an EDGE_ABNORMAL edge from entry to exit.

	* testsuite/libgomp.c/pr66133.c: New test.

From-SVN: r223171
2015-05-13 18:59:20 +02:00
Michael Haubenwallner 55fba60115 libgomp: Bump to automake 1.11.6
2015-05-13  Michael Haubenwallner  <michael.haubenwallner@ssi-schaefer.com>

	* Makefile.in: Regenerated with automake-1.11.6.
	* aclocal.m4: Likewise.
	* config.h.in: Likewise.
	* configure: Likewise.
	* testsuite/Makefile.in: Likewise.

From-SVN: r223143
2015-05-13 11:24:38 +00:00
Jason Merrill 1be0899d94 openacc.h (__GOACC_NOTHROW): Fix noexcept syntax.
* openacc.h (__GOACC_NOTHROW): Fix noexcept syntax.

	* testsuite/libgomp.oacc-c-c++-common/reduction-4.c (main): Use
	_Complex.

From-SVN: r222966
2015-05-09 00:50:35 -04:00
Thomas Schwinge ae8ffbbb8d [PR testsuite/65205, libgomp/65993] Fix dg-shouldfail usage in OpenACC libgomp tests
In dg-output, don't expect "0x" prefix for "%p" format specifier, don't expect
"(nil)" for NULL pointer.

	PR testsuite/65205
	PR libgomp/65993
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/clauses-2.c: In dg-output,
	don't expect "0x" prefix for "%p" format specifier, don't expect
	"(nil)" for NULL pointer.
	* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-26.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-27.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-35.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-36.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-39.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-40.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-57.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-58.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: More
	accurately specify what we're looking for.
	* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.

From-SVN: r222799
2015-05-05 11:39:29 +02:00
James Norris 1309f1d25b [PR testsuite/65205] Fix dg-shouldfail usage in OpenACC libgomp tests
PR testsuite/65205
	libgomp/
	* testsuite/lib/libgomp.exp
	(check_effective_target_openacc_host_selected)
	(check_effective_target_openacc_host_nonshm_selected): New
	procedures.
	* testsuite/libgomp.oacc-c-c++-common/clauses-2.c: Fix misuse of
	dg-shouldfail.
	* testsuite/libgomp.oacc-c-c++-common/lib-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-11.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-26.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-27.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-35.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-36.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-39.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-40.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-57.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-58.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-62.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-63.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-64.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-65.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-67.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-68.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-71.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-77.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-80.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.

From-SVN: r222620
2015-04-30 14:44:39 +02:00
Julian Brown d93bdab53b mkoffload.c (process): Support variable mapping.
gcc/
	* config/nvptx/mkoffload.c (process): Support variable mapping.

	libgomp/
	* libgomp.h (target_mem_desc: Remove mem_map field.
	(acc_dispatch_t): Remove open_device_func, close_device_func,
	get_device_num_func, set_device_num_func, target_data members.
	Change create_thread_data_func argument to device number instead of
	generic pointer.
	* oacc-async.c (assert.h): Include.
	(acc_async_test, acc_async_test_all, acc_wait, acc_wait_async)
	(acc_wait_all, acc_wait_all_async): Use current host thread's
	active device, not base_dev.
	* oacc-cuda.c (acc_get_current_cuda_device)
	(acc_get_current_cuda_context, acc_get_cuda_stream)
	(acc_set_cuda_stream): Likewise.
	* oacc-host.c (host_dispatch): Don't set open_device_func,
	close_device_func, get_device_num_func or set_device_num_func.
	* oacc-init.c (base_dev, init_key): Remove.
	(cached_base_dev): New.
	(name_of_acc_device_t): New.
	(acc_init_1): Initialise default-numbered device, not zeroth.
	(acc_shutdown_1): Close all devices of a given type.
	(goacc_destroy_thread): Don't use base_dev.
	(lazy_open, lazy_init, lazy_init_and_open): Remove.
	(goacc_attach_host_thread_to_device): New.
	(acc_init): Reimplement with goacc_attach_host_thread_to_device.
	(acc_get_num_devices): Don't use base_dev.
	(acc_set_device_type): Reimplement.
	(acc_get_device_type): Don't use base_dev.
	(acc_get_device_num): Tweak logic.
	(acc_set_device_num): Likewise.
	(acc_on_device): Use acc_get_device_type.
	(goacc_runtime_initialize): Initialize cached_base_dev not base_dev.
	(goacc_lazy_initialize): Reimplement with acc_init and
	goacc_attach_host_thread_to_device.
	* oacc-int.h (goacc_thread): Add base_dev field.
	(base_dev): Remove extern declaration.
	(goacc_attach_host_thread_to_device): Add prototype.
	* oacc-mem.c (acc_malloc): Use current thread's device instead of
	base_dev.
	(acc_free): Likewise.
	(acc_memcpy_to_device): Likewise.
	(acc_memcpy_from_device): Likewise.
	* oacc-parallel.c (select_acc_device): Remove. Replace calls with
	goacc_lazy_initialize (throughout).
	(GOACC_parallel): Use tgt_offset to locate target functions.
	* target.c (gomp_map_vars): Don't set tgt->mem_map.
	(gomp_unmap_vars): Use devicep->mem_map pointer not tgt->mem_map.
	(gomp_load_plugin_for_device): Remove open_device, close_device,
	get_device_num, set_device_num openacc hook initialisation. Don't set
	openacc.target_data.
	* plugin/plugin-host.c (GOMP_OFFLOAD_openacc_open_device)
	(GOMP_OFFLOAD_openacc_close_device)
	(GOMP_OFFLOAD_openacc_get_device_num)
	(GOMP_OFFLOAD_openacc_set_device_num): Remove.
	(GOMP_OFFLOAD_openacc_create_thread_data): Change (unused) argument
	to int.
	* plugin/plugin-nvptx.c (ptx_inited): Remove.
	(instantiated_devices, ptx_dev_lock): New.
	(struct ptx_image_data): New.
	(ptx_devices, ptx_images, ptx_image_lock): New.
	(fini_streams_for_device): Reorder cuStreamDestroy call.
	(nvptx_get_num_devices): Remove forward declaration.
	(nvptx_init): Change return type to bool.
	(nvptx_fini): Remove.
	(nvptx_attach_host_thread_to_device): New.
	(nvptx_open_device): Return struct ptx_device* instead of void*.
	(nvptx_close_device): Change argument type to struct ptx_device*,
	return type to void.
	(nvptx_get_num_devices): Use instantiated_devices not ptx_inited.
	(kernel_target_data, kernel_host_table): Remove static globals.
	(GOMP_OFFLOAD_register_image, GOMP_OFFLOAD_get_table): Remove.
	(GOMP_OFFLOAD_init_device): Reimplement.
	(GOMP_OFFLOAD_fini_device): Likewise.
	(GOMP_OFFLOAD_load_image, GOMP_OFFLOAD_unload_image): New.
	(GOMP_OFFLOAD_alloc, GOMP_OFFLOAD_free, GOMP_OFFLOAD_dev2host)
	(GOMP_OFFLOAD_host2dev): Use ORD argument.
	(GOMP_OFFLOAD_openacc_open_device)
	(GOMP_OFFLOAD_openacc_close_device)
	(GOMP_OFFLOAD_openacc_set_device_num)
	(GOMP_OFFLOAD_openacc_get_device_num): Remove.
	(GOMP_OFFLOAD_openacc_create_thread_data): Change argument to int
	(device number).

	libgomp/testsuite/
	* libgomp.oacc-c-c++-common/lib-9.c: Fix devnum check in test.

From-SVN: r221922
2015-04-08 15:58:33 +00:00
Jakub Jelinek 6c384511a3 re PR middle-end/65597 (ICE in build_outer_var_ref, at omp-low.c:1043)
PR fortran/65597
	* trans-openmp.c (gfc_trans_omp_do): For !simple simd with explicit
	linear clause for the iterator set OMP_CLAUSE_LINEAR_NO_COPYIN.
	For implcitly added !simple OMP_CLAUSE_LINEAR set it too.  Use step 1
	instead of the original step on the new iterator - count.

	* testsuite/libgomp.fortran/pr65597.f90: New test.

From-SVN: r221776
2015-03-30 19:54:05 +02:00
Tom de Vries 3e8165a5df Add verification to libgomp.graphite/force-parallel-6.c
2015-03-27  Tom de Vries  <tom@codesourcery.com>

	PR testsuite/65594
	* testsuite/libgomp.graphite/force-parallel-6.c (abort): Declare.
	(init, check): New function.
	(foo): Change return type to void.
	(main): Call init and check.

From-SVN: r221728
2015-03-27 12:10:16 +00:00
Tom de Vries 4d688c9a17 Scale down libgomp.graphite/force-parallel-6.c
2015-03-27  Tom de Vries  <tom@codesourcery.com>

	PR testsuite/65594
	* testsuite/libgomp.graphite/force-parallel-6.c (M): Define.
	(foo): Use M for non-inner loops to scale down test-case.

From-SVN: r221727
2015-03-27 12:10:07 +00:00
John David Anglin 844d9a76bc reduction-4.c: Don't run on hppa*-*-hpux*.
* testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Don't run on
	hppa*-*-hpux*.

From-SVN: r221557
2015-03-21 16:19:43 +00:00
Jakub Jelinek 9b65e171ed c-decl.c (c_decl_attributes): Also add "omp declare target" attribute for DECL_EXTERNAL VAR_DECLs.
* c-decl.c (c_decl_attributes): Also add "omp declare target"
	attribute for DECL_EXTERNAL VAR_DECLs.

	* decl2.c (cplus_decl_attributes): Also add "omp declare target"
	attribute for DECL_EXTERNAL VAR_DECLs.

	* testsuite/libgomp.c/target-10.c: New test.
	* testsuite/libgomp.c++/target-4.C: New test.

From-SVN: r221520
2015-03-19 20:12:43 +01:00
Ilya Verbin db397e2e5a varpool.c (varpool_node::get_create): Don't set 'offloadable' flag for the external decls.
gcc/
	* varpool.c (varpool_node::get_create): Don't set 'offloadable' flag for
	the external decls.
libgomp/
	* testsuite/libgomp.fortran/declare-target-1.f90: New test.
	* testsuite/libgomp.fortran/declare-target-2.f90: New file.

From-SVN: r221421
2015-03-13 13:30:26 +00:00
Tom de Vries 05deba9fe0 Use DO_PRAGMA in libgomp.oacc-c-c++-common/reduction-1.c
2015-02-25  Tom de Vries  <tom@codesourcery.com>

	* testsuite/libgomp.oacc-c-c++-common/reduction-1.c (DO_PRAGMA)
	(check_reduction_op, check_reduction_macro, max, min):
	Declare.
	(test_reductions_int, test_reductions_minmax, test_reductions_bool): New
	function.
	(main): Use new functions.

From-SVN: r220971
2015-02-25 15:38:38 +00:00
Jakub Jelinek fa01ffccc7 re PR middle-end/64824 (ICE in gimple verification)
PR c/64824
	* c-parser.c (c_parser_binary_expression): Fix OpenMP stack[sp].prec
	check in the POP macro.

	* testsuite/libgomp.c/atomic-18.c: New test.
	* testsuite/libgomp.c++/atomic-16.C: New test.

From-SVN: r220617
2015-02-11 14:40:31 +01:00
Jakub Jelinek 4886ec8e70 re PR middle-end/64824 (ICE in gimple verification)
PR c/64824
	PR c/64868
gcc/c/
	* c-parser.c (c_parser_omp_atomic): Handle RDIV_EXPR.
gcc/cp/
	* parser.c (cp_parser_omp_atomic): Handle RDIV_EXPR.
gcc/c-family/
	* c-omp.c (c_finish_omp_atomic): Use TRUNC_DIV_EXPR
	instead of RDIV_EXPR.  Use build_binary_op instead of
	build2_loc.
libgomp/
	* testsuite/libgomp.c/pr64824.c: New test.
	* testsuite/libgomp.c/pr64868.c: New test.
	* testsuite/libgomp.c++/pr64824.C: New test.
	* testsuite/libgomp.c++/pr64868.C: New test.

From-SVN: r220420
2015-02-04 23:32:36 +01:00
Jakub Jelinek b891044723 re PR middle-end/64734 (ICE at omp lowering)
PR middle-end/64734
	* omp-low.c (scan_sharing_clauses): Don't ignore
	OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION GOMP_MAP_POINTER clauses
	on target data/update constructs.

	* libgomp.c/pr64734.c: New test.

From-SVN: r220053
2015-01-23 19:19:50 +01:00
Tom de Vries a0c88d0629 Make fopenacc an LTO option
2015-01-23  Tom de Vries  <tom@codesourcery.com>

	PR libgomp/64672
	* lto-opts.c (lto_write_options): Output non-explicit conservative
	-fno-openacc.
	* lto-wrapper.c (merge_and_complain): Handle merging -fopenacc.
	(append_compiler_options): Pass -fopenacc through.

	* c.opt (fopenacc): Mark as LTO option.

	* lang.opt (fopenacc): Mark as LTO option.

	* testsuite/libgomp.oacc-c-c++-common/abort-5.c: New test.

From-SVN: r220038
2015-01-23 12:54:16 +00:00
Tom de Vries 1506ae0e1e Make fopenmp an LTO option
2015-01-23  Tom de Vries  <tom@codesourcery.com>

	PR libgomp/64707
	* lto-opts.c (lto_write_options): Output non-explicit conservative
	-fno-openmp.
	* lto-wrapper.c (merge_and_complain): Handle merging -fopenmp.
	(append_compiler_options): Pass -fopenmp through.

	* c.opt (fopenmp): Mark as LTO option.

	* lang.opt (fopenmp): Mark as LTO option.

	* testsuite/libgomp.c/target-9.c: Add -ftree-parallelize-loops=0 to
	dg-options.

From-SVN: r220037
2015-01-23 12:53:55 +00:00
Thomas Schwinge 41dbbb3789 Merge current set of OpenACC changes from gomp-4_0-branch.
contrib/
	* gcc_update (files_and_dependencies): Update rules for new
	libgomp/plugin/Makefrag.am and libgomp/plugin/configfrag.ac files.
	gcc/
	* builtin-types.def (BT_FN_VOID_INT_INT_VAR)
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
	New function types.
	* builtins.c: Include "gomp-constants.h".
	(expand_builtin_acc_on_device): New function.
	(expand_builtin, is_inexpensive_builtin): Handle
	BUILT_IN_ACC_ON_DEVICE.
	* builtins.def (DEF_GOACC_BUILTIN, DEF_GOACC_BUILTIN_COMPILER):
	New macros.
	* cgraph.c (cgraph_node::create): Consider flag_openacc next to
	flag_openmp.
	* config.gcc <nvptx-*> (tm_file): Add nvptx/offload.h.
	<*-intelmic-* | *-intelmicemul-*> (tm_file): Add
	i386/intelmic-offload.h.
	* gcc.c (LINK_COMMAND_SPEC, GOMP_SELF_SPECS): For -fopenacc, link
	to libgomp and its dependencies.
	* config/arc/arc.h (LINK_COMMAND_SPEC): Likewise.
	* config/darwin.h (LINK_COMMAND_SPEC_A): Likewise.
	* config/i386/mingw32.h (GOMP_SELF_SPECS): Likewise.
	* config/ia64/hpux.h (LIB_SPEC): Likewise.
	* config/pa/pa-hpux11.h (LIB_SPEC): Likewise.
	* config/pa/pa64-hpux.h (LIB_SPEC): Likewise.
	* doc/generic.texi: Update for OpenACC changes.
	* doc/gimple.texi: Likewise.
	* doc/invoke.texi: Likewise.
	* doc/sourcebuild.texi: Likewise.
	* gimple-pretty-print.c (dump_gimple_omp_for): Handle
	GF_OMP_FOR_KIND_OACC_LOOP.
	(dump_gimple_omp_target): Handle GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_DATA,
	GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
	Dump more data.
	* gimple.c: Update comments for OpenACC changes.
	* gimple.def: Likewise.
	* gimple.h: Likewise.
	(enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP,
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_DATA, GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
	(gimple_omp_for_cond, gimple_omp_for_set_cond): Sort in the
	appropriate place.
	(is_gimple_omp_oacc, is_gimple_omp_offloaded): New functions.
	* gimplify.c: Include "gomp-constants.h".
	Update comments for OpenACC changes.
	(is_gimple_stmt): Handle OACC_PARALLEL, OACC_KERNELS, OACC_DATA,
	OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE, OACC_ENTER_DATA,
	OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
	(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
	OMP_CLAUSE__CACHE_, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER,
	OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
	OMP_CLAUSE_SEQ.
	(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(gimplify_oacc_cache): New function.
	(gimplify_omp_for): Handle OACC_LOOP.
	(gimplify_omp_workshare): Handle OACC_KERNELS, OACC_PARALLEL,
	OACC_DATA.
	(gimplify_omp_target_update): Handle OACC_ENTER_DATA,
	OACC_EXIT_DATA, OACC_UPDATE.
	(gimplify_expr): Handle OACC_LOOP, OACC_CACHE, OACC_HOST_DATA,
	OACC_DECLARE, OACC_KERNELS, OACC_PARALLEL, OACC_DATA,
	OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE.
	(gimplify_body): Consider flag_openacc next to flag_openmp.
	* lto-streamer-out.c: Include "gomp-constants.h".
	* omp-builtins.def (BUILT_IN_ACC_GET_DEVICE_TYPE)
	(BUILT_IN_GOACC_DATA_START, BUILT_IN_GOACC_DATA_END)
	(BUILT_IN_GOACC_ENTER_EXIT_DATA, BUILT_IN_GOACC_PARALLEL)
	(BUILT_IN_GOACC_UPDATE, BUILT_IN_GOACC_WAIT)
	(BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS)
	(BUILT_IN_ACC_ON_DEVICE): New builtins.
	* omp-low.c: Include "gomp-constants.h".
	Update comments for OpenACC changes.
	(struct omp_context): Add reduction_map, gwv_below, gwv_this
	members.
	(extract_omp_for_data, use_pointer_for_field, install_var_field)
	(new_omp_context, delete_omp_context, scan_sharing_clauses)
	(create_omp_child_function, scan_omp_for, scan_omp_target)
	(check_omp_nesting_restrictions, lower_reduction_clauses)
	(build_omp_regions_1, diagnose_sb_0, make_gimple_omp_edges):
	Update for OpenACC changes.
	(scan_sharing_clauses): Handle OMP_CLAUSE_NUM_GANGS:
	OMP_CLAUSE_NUM_WORKERS: OMP_CLAUSE_VECTOR_LENGTH,
	OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT, OMP_CLAUSE_GANG,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_INDEPENDENT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ.  Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.
	(expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
	Handle GF_OMP_FOR_KIND_OACC_LOOP.
	(expand_omp_target, lower_omp_target): Handle
	GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
	GF_OMP_TARGET_KIND_OACC_UPDATE,
	GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA,
	GF_OMP_TARGET_KIND_OACC_DATA.
	(pass_expand_omp::execute, execute_lower_omp)
	(pass_diagnose_omp_blocks::gate): Consider flag_openacc next to
	flag_openmp.
	(offload_symbol_decl): New variable.
	(oacc_get_reduction_array_id, oacc_max_threads)
	(get_offload_symbol_decl, get_base_type, lookup_oacc_reduction)
	(maybe_lookup_oacc_reduction, enclosing_target_ctx)
	(oacc_loop_or_target_p, oacc_lower_reduction_var_helper)
	(oacc_gimple_assign, oacc_initialize_reduction_data)
	(oacc_finalize_reduction_data, oacc_process_reduction_data): New
	functions.
	(is_targetreg_ctx): Remove function.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CACHE_,
	OMP_CLAUSE_DEVICE_RESIDENT, OMP_CLAUSE_USE_DEVICE,
	OMP_CLAUSE_GANG, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_NUM_GANGS,
	OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH.
	* tree.c (omp_clause_code_name, walk_tree_1): Update accordingly.
	* tree.h (OMP_CLAUSE_GANG_EXPR, OMP_CLAUSE_GANG_STATIC_EXPR)
	(OMP_CLAUSE_ASYNC_EXPR, OMP_CLAUSE_WAIT_EXPR)
	(OMP_CLAUSE_VECTOR_EXPR, OMP_CLAUSE_WORKER_EXPR)
	(OMP_CLAUSE_NUM_GANGS_EXPR, OMP_CLAUSE_NUM_WORKERS_EXPR)
	(OMP_CLAUSE_VECTOR_LENGTH_EXPR): New macros.
	* tree-core.h: Update comments for OpenACC changes.
	(enum omp_clause_map_kind): Remove.
	(struct tree_omp_clause): Change type of map_kind member from enum
	omp_clause_map_kind to unsigned char.
	* tree-inline.c: Update comments for OpenACC changes.
	* tree-nested.c: Likewise.  Include "gomp-constants.h".
	(convert_nonlocal_reference_stmt, convert_local_reference_stmt)
	(convert_tramp_reference_stmt, convert_gimple_call): Update for
	OpenACC changes.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	* tree-pretty-print.c: Include "gomp-constants.h".
	(dump_omp_clause): Handle OMP_CLAUSE_DEVICE_RESIDENT,
	OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_GANG,
	OMP_CLAUSE_ASYNC, OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ,
	OMP_CLAUSE_WAIT, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_INDEPENDENT.  Use GOMP_MAP_*
	instead of OMP_CLAUSE_MAP_*.
	(dump_generic_node): Handle OACC_PARALLEL, OACC_KERNELS,
	OACC_DATA, OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE,
	OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
	* tree-streamer-in.c: Include "gomp-constants.h".
	(unpack_ts_omp_clause_value_fields) Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.  Use OMP_CLAUSE_SET_MAP_KIND.
	* tree-streamer-out.c: Include "gomp-constants.h".
	(pack_ts_omp_clause_value_fields): Use GOMP_MAP_* instead of
	OMP_CLAUSE_MAP_*.
	* tree.def (OACC_PARALLEL, OACC_KERNELS, OACC_DATA)
	(OACC_HOST_DATA, OACC_LOOP, OACC_CACHE, OACC_DECLARE)
	(OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE): New tree codes.
	* tree.c (omp_clause_num_ops): Update accordingly.
	* tree.h (OMP_BODY, OMP_CLAUSES, OMP_LOOP_CHECK, OMP_CLAUSE_SIZE):
	Likewise.
	(OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES, OACC_KERNELS_BODY)
	(OACC_KERNELS_CLAUSES, OACC_DATA_BODY, OACC_DATA_CLAUSES)
	(OACC_HOST_DATA_BODY, OACC_HOST_DATA_CLAUSES, OACC_CACHE_CLAUSES)
	(OACC_DECLARE_CLAUSES, OACC_ENTER_DATA_CLAUSES)
	(OACC_EXIT_DATA_CLAUSES, OACC_UPDATE_CLAUSES)
	(OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): New macros.
	* tree.h (OMP_CLAUSE_MAP_KIND): Cast it to enum gomp_map_kind.
	(OMP_CLAUSE_SET_MAP_KIND): New macro.
	* varpool.c (varpool_node::get_create): Consider flag_openacc next
	to flag_openmp.
	* config/i386/intelmic-offload.h: New file.
	* config/nvptx/offload.h: Likewise.
	gcc/ada/
	* gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_8)
	(DEF_FUNCTION_TYPE_VAR_12): New macros.
	gcc/c-family/
	* c.opt (fopenacc): New option.
	* c-cppbuiltin.c (c_cpp_builtins): Conditionally define _OPENACC.
	* c-common.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
	New macros.
	* c-common.h (c_finish_oacc_wait): New prototype.
	* c-omp.c: Include "omp-low.h" and "gomp-constants.h".
	(c_finish_oacc_wait): New function.
	* c-pragma.c (oacc_pragmas): New variable.
	(c_pp_lookup_pragma, init_pragma): Handle it.
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_CACHE,
	PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA,
	PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL,
	PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
	(enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ASYNC,
	PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_COLLAPSE,
	PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYIN,
	PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE,
	PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICE,
	PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE,
	PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST,
	PRAGMA_OACC_CLAUSE_IF, PRAGMA_OACC_CLAUSE_NUM_GANGS,
	PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
	PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_PRIVATE,
	PRAGMA_OACC_CLAUSE_REDUCTION, PRAGMA_OACC_CLAUSE_SELF,
	PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_VECTOR,
	PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT,
	PRAGMA_OACC_CLAUSE_WORKER.
	gcc/c/
	* c-parser.c: Include "gomp-constants.h".
	(c_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
	omp_clause_map_kind.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	Use OMP_CLAUSE_SET_MAP_KIND.
	(c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA,
	PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_UPDATE.
	(c_parser_omp_construct): Handle PRAGMA_OACC_CACHE,
	PRAGMA_OACC_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP,
	PRAGMA_OACC_PARALLEL, PRAGMA_OACC_WAIT.
	(c_parser_omp_clause_name): Handle "auto", "async", "copy",
	"copyout", "create", "delete", "deviceptr", "gang", "host",
	"num_gangs", "num_workers", "present", "present_or_copy", "pcopy",
	"present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
	"present_or_create", "pcreate", "seq", "self", "vector",
	"vector_length", "wait", "worker".
	(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
	(OACC_ENTER_DATA_CLAUSE_MASK, OACC_EXIT_DATA_CLAUSE_MASK)
	(OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
	(OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
	(c_parser_omp_variable_list): Handle OMP_CLAUSE__CACHE_.
	(c_parser_oacc_wait_list, c_parser_oacc_data_clause)
	(c_parser_oacc_data_clause_deviceptr)
	(c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_workers)
	(c_parser_oacc_clause_async, c_parser_oacc_clause_wait)
	(c_parser_omp_clause_vector_length, c_parser_oacc_all_clauses)
	(c_parser_oacc_cache, c_parser_oacc_data, c_parser_oacc_kernels)
	(c_parser_oacc_enter_exit_data, c_parser_oacc_loop)
	(c_parser_oacc_parallel, c_parser_oacc_update)
	(c_parser_oacc_wait): New functions.
	* c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels)
	(c_finish_oacc_data): New prototypes.
	* c-typeck.c: Include "gomp-constants.h".
	(handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR.  Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(c_finish_oacc_parallel, c_finish_oacc_kernels)
	(c_finish_oacc_data): New functions.
	(c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_,
	OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
	OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_GANG,
	OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, and OMP_CLAUSE_MAP's
	GOMP_MAP_FORCE_DEVICEPTR.
	gcc/cp/
	* parser.c: Include "gomp-constants.h".
	(cp_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
	omp_clause_map_kind.  Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	Use OMP_CLAUSE_SET_MAP_KIND.
	(cp_parser_omp_construct, cp_parser_pragma): Handle
	PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA,
	PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_PARALLEL,
	PRAGMA_OACC_LOOP, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
	(cp_parser_omp_clause_name): Handle "async", "copy", "copyout",
	"create", "delete", "deviceptr", "host", "num_gangs",
	"num_workers", "present", "present_or_copy", "pcopy",
	"present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
	"present_or_create", "pcreate", "vector_length", "wait".
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
	(OACC_EXIT_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
	(OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
	(OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
	(cp_parser_omp_var_list_no_open): Handle OMP_CLAUSE__CACHE_.
	(cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr)
	(cp_parser_oacc_clause_vector_length, cp_parser_oacc_wait_list)
	(cp_parser_oacc_clause_wait, cp_parser_omp_clause_num_gangs)
	(cp_parser_omp_clause_num_workers, cp_parser_oacc_clause_async)
	(cp_parser_oacc_all_clauses, cp_parser_oacc_cache)
	(cp_parser_oacc_data, cp_parser_oacc_enter_exit_data)
	(cp_parser_oacc_kernels, cp_parser_oacc_loop)
	(cp_parser_oacc_parallel, cp_parser_oacc_update)
	(cp_parser_oacc_wait): New functions.
	* cp-tree.h (finish_oacc_data, finish_oacc_kernels)
	(finish_oacc_parallel): New prototypes.
	* semantics.c: Include "gomp-constants.h".
	(handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR.  Use
	GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.  Use
	OMP_CLAUSE_SET_MAP_KIND.
	(finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
	OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT, OMP_CLAUSE__CACHE_.
	Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
	(finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
	functions.
	gcc/fortran/
	* lang.opt (fopenacc): New option.
	* cpp.c (cpp_define_builtins): Conditionally define _OPENACC.
	* dump-parse-tree.c (show_omp_node): Split part of it into...
	(show_omp_clauses): ... this new function.
	(show_omp_node, show_code_node): Handle EXEC_OACC_PARALLEL_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	(show_namespace): Update for OpenACC.
	* f95-lang.c (DEF_FUNCTION_TYPE_VAR_2, DEF_FUNCTION_TYPE_VAR_8)
	(DEF_FUNCTION_TYPE_VAR_12, DEF_GOACC_BUILTIN)
	(DEF_GOACC_BUILTIN_COMPILER): New macros.
	* types.def (BT_FN_VOID_INT_INT_VAR)
	(BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
	New function types.
	* gfortran.h (gfc_statement): Add ST_OACC_PARALLEL_LOOP,
	ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_DATA,
	ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA,
	ST_OACC_LOOP, ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE,
	ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP,
	ST_OACC_END_KERNELS_LOOP, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
	ST_OACC_ROUTINE.
	(struct gfc_expr_list): New data type.
	(gfc_get_expr_list): New macro.
	(gfc_omp_map_op): Add OMP_MAP_FORCE_ALLOC, OMP_MAP_FORCE_DEALLOC,
	OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM, OMP_MAP_FORCE_TOFROM,
	OMP_MAP_FORCE_PRESENT, OMP_MAP_FORCE_DEVICEPTR.
	(OMP_LIST_FIRST, OMP_LIST_DEVICE_RESIDENT, OMP_LIST_USE_DEVICE)
	(OMP_LIST_CACHE): New enumerators.
	(struct gfc_omp_clauses): Add async_expr, gang_expr, worker_expr,
	vector_expr, num_gangs_expr, num_workers_expr, vector_length_expr,
	wait_list, tile_list, async, gang, worker, vector, seq,
	independent, wait, par_auto, gang_static, and loc members.
	(struct gfc_namespace): Add oacc_declare_clauses member.
	(gfc_exec_op): Add EXEC_OACC_KERNELS_LOOP,
	EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	(gfc_free_expr_list, gfc_resolve_oacc_directive)
	(gfc_resolve_oacc_declare, gfc_resolve_oacc_parallel_loop_blocks)
	(gfc_resolve_oacc_blocks): New prototypes.
	* match.c (match_exit_cycle): Handle EXEC_OACC_LOOP and
	EXEC_OACC_PARALLEL_LOOP.
	* match.h (gfc_match_oacc_cache, gfc_match_oacc_wait)
	(gfc_match_oacc_update, gfc_match_oacc_declare)
	(gfc_match_oacc_loop, gfc_match_oacc_host_data)
	(gfc_match_oacc_data, gfc_match_oacc_kernels)
	(gfc_match_oacc_kernels_loop, gfc_match_oacc_parallel)
	(gfc_match_oacc_parallel_loop, gfc_match_oacc_enter_data)
	(gfc_match_oacc_exit_data, gfc_match_oacc_routine): New
	prototypes.
	* openmp.c: Include "diagnostic.h" and "gomp-constants.h".
	(gfc_free_omp_clauses): Update for members added to struct
	gfc_omp_clauses.
	(gfc_match_omp_clauses): Change mask paramter to uint64_t.  Add
	openacc parameter.
	(resolve_omp_clauses): Add openacc parameter.  Update for OpenACC.
	(struct fortran_omp_context): Add is_openmp member.
	(gfc_resolve_omp_parallel_blocks): Initialize it.
	(gfc_resolve_do_iterator): Update for OpenACC.
	(gfc_resolve_omp_directive): Call
	resolve_omp_directive_inside_oacc_region.
	(OMP_CLAUSE_PRIVATE, OMP_CLAUSE_FIRSTPRIVATE)
	(OMP_CLAUSE_LASTPRIVATE, OMP_CLAUSE_COPYPRIVATE)
	(OMP_CLAUSE_SHARED, OMP_CLAUSE_COPYIN, OMP_CLAUSE_REDUCTION)
	(OMP_CLAUSE_IF, OMP_CLAUSE_NUM_THREADS, OMP_CLAUSE_SCHEDULE)
	(OMP_CLAUSE_DEFAULT, OMP_CLAUSE_ORDERED, OMP_CLAUSE_COLLAPSE)
	(OMP_CLAUSE_UNTIED, OMP_CLAUSE_FINAL, OMP_CLAUSE_MERGEABLE)
	(OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND, OMP_CLAUSE_INBRANCH)
	(OMP_CLAUSE_LINEAR, OMP_CLAUSE_NOTINBRANCH, OMP_CLAUSE_PROC_BIND)
	(OMP_CLAUSE_SAFELEN, OMP_CLAUSE_SIMDLEN, OMP_CLAUSE_UNIFORM)
	(OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO)
	(OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT)
	(OMP_CLAUSE_DIST_SCHEDULE): Use uint64_t.
	(OMP_CLAUSE_ASYNC, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS)
	(OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_COPY, OMP_CLAUSE_COPYOUT)
	(OMP_CLAUSE_CREATE, OMP_CLAUSE_PRESENT)
	(OMP_CLAUSE_PRESENT_OR_COPY, OMP_CLAUSE_PRESENT_OR_COPYIN)
	(OMP_CLAUSE_PRESENT_OR_COPYOUT, OMP_CLAUSE_PRESENT_OR_CREATE)
	(OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER)
	(OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT)
	(OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_DEVICE_RESIDENT)
	(OMP_CLAUSE_HOST_SELF, OMP_CLAUSE_OACC_DEVICE, OMP_CLAUSE_WAIT)
	(OMP_CLAUSE_DELETE, OMP_CLAUSE_AUTO, OMP_CLAUSE_TILE): New macros.
	(gfc_match_omp_clauses): Handle those.
	(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
	(OACC_LOOP_CLAUSES, OACC_PARALLEL_LOOP_CLAUSES)
	(OACC_KERNELS_LOOP_CLAUSES, OACC_HOST_DATA_CLAUSES)
	(OACC_DECLARE_CLAUSES, OACC_UPDATE_CLAUSES)
	(OACC_ENTER_DATA_CLAUSES, OACC_EXIT_DATA_CLAUSES)
	(OACC_WAIT_CLAUSES): New macros.
	(gfc_free_expr_list, match_oacc_expr_list, match_oacc_clause_gang)
	(gfc_match_omp_map_clause, gfc_match_oacc_parallel_loop)
	(gfc_match_oacc_parallel, gfc_match_oacc_kernels_loop)
	(gfc_match_oacc_kernels, gfc_match_oacc_data)
	(gfc_match_oacc_host_data, gfc_match_oacc_loop)
	(gfc_match_oacc_declare, gfc_match_oacc_update)
	(gfc_match_oacc_enter_data, gfc_match_oacc_exit_data)
	(gfc_match_oacc_wait, gfc_match_oacc_cache)
	(gfc_match_oacc_routine, oacc_is_loop)
	(resolve_oacc_scalar_int_expr, resolve_oacc_positive_int_expr)
	(check_symbol_not_pointer, check_array_not_assumed)
	(resolve_oacc_data_clauses, resolve_oacc_deviceptr_clause)
	(oacc_compatible_clauses, oacc_is_parallel, oacc_is_kernels)
	(omp_code_to_statement, oacc_code_to_statement)
	(resolve_oacc_directive_inside_omp_region)
	(resolve_omp_directive_inside_oacc_region)
	(resolve_oacc_nested_loops, resolve_oacc_params_in_parallel)
	(resolve_oacc_loop_blocks, gfc_resolve_oacc_blocks)
	(resolve_oacc_loop, resolve_oacc_cache, gfc_resolve_oacc_declare)
	(gfc_resolve_oacc_directive): New functions.
	* parse.c (next_free): Update for OpenACC.  Move some code into...
	(verify_token_free): ... this new function.
	(next_fixed): Update for OpenACC.  Move some code into...
	(verify_token_fixed): ... this new function.
	(case_executable): Add ST_OACC_UPDATE, ST_OACC_WAIT,
	ST_OACC_CACHE, ST_OACC_ENTER_DATA, and ST_OACC_EXIT_DATA.
	(case_exec_markers): Add ST_OACC_PARALLEL_LOOP, ST_OACC_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA, ST_OACC_LOOP,
	ST_OACC_KERNELS_LOOP.
	(case_decl): Add ST_OACC_ROUTINE.
	(push_state, parse_critical_block, parse_progunit): Update for
	OpenACC.
	(gfc_ascii_statement): Handle ST_OACC_PARALLEL_LOOP,
	ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_KERNELS_LOOP,
	ST_OACC_END_KERNELS_LOOP, ST_OACC_DATA, ST_OACC_END_DATA,
	ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP,
	ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
	ST_OACC_CACHE, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
	ST_OACC_ROUTINE.
	(verify_st_order, parse_spec): Handle ST_OACC_DECLARE.
	(parse_executable): Handle ST_OACC_PARALLEL_LOOP,
	ST_OACC_KERNELS_LOOP, ST_OACC_LOOP, ST_OACC_PARALLEL,
	ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA.
	(decode_oacc_directive, parse_oacc_structured_block)
	(parse_oacc_loop, is_oacc): New functions.
	* parse.h (struct gfc_state_data): Add oacc_declare_clauses
	member.
	(is_oacc): New prototype.
	* resolve.c (gfc_resolve_blocks, gfc_resolve_code): Handle
	EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL,
	EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS, EXEC_OACC_DATA,
	EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE,
	EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA,
	EXEC_OACC_EXIT_DATA.
	(resolve_codes): Call gfc_resolve_oacc_declare.
	* scanner.c (openacc_flag, openacc_locus): New variables.
	(skip_free_comments): Update for OpenACC.  Move some code into...
	(skip_omp_attribute): ... this new function.
	(skip_oacc_attribute): New function.
	(skip_fixed_comments, gfc_next_char_literal): Update for OpenACC.
	* st.c (gfc_free_statement): Handle EXEC_OACC_PARALLEL_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
	EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
	EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
	* trans-decl.c (gfc_generate_function_code): Update for OpenACC.
	* trans-openmp.c: Include "gomp-constants.h".
	(gfc_omp_finish_clause, gfc_trans_omp_clauses): Use GOMP_MAP_*
	instead of OMP_CLAUSE_MAP_*.  Use OMP_CLAUSE_SET_MAP_KIND.
	(gfc_trans_omp_clauses): Handle OMP_LIST_USE_DEVICE,
	OMP_LIST_DEVICE_RESIDENT, OMP_LIST_CACHE, and OMP_MAP_FORCE_ALLOC,
	OMP_MAP_FORCE_DEALLOC, OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM,
	OMP_MAP_FORCE_TOFROM, OMP_MAP_FORCE_PRESENT,
	OMP_MAP_FORCE_DEVICEPTR, and gfc_omp_clauses' async, seq,
	independent, wait_list, num_gangs_expr, num_workers_expr,
	vector_length_expr, vector, vector_expr, worker, worker_expr,
	gang, gang_expr members.
	(gfc_trans_omp_do): Handle EXEC_OACC_LOOP.
	(gfc_convert_expr_to_tree, gfc_trans_oacc_construct)
	(gfc_trans_oacc_executable_directive)
	(gfc_trans_oacc_wait_directive, gfc_trans_oacc_combined_directive)
	(gfc_trans_oacc_declare, gfc_trans_oacc_directive): New functions.
	* trans-stmt.c (gfc_trans_block_construct): Update for OpenACC.
	* trans-stmt.h (gfc_trans_oacc_directive, gfc_trans_oacc_declare):
	New prototypes.
	* trans.c (tranc_code): Handle EXEC_OACC_CACHE, EXEC_OACC_WAIT,
	EXEC_OACC_UPDATE, EXEC_OACC_LOOP, EXEC_OACC_HOST_DATA,
	EXEC_OACC_DATA, EXEC_OACC_KERNELS, EXEC_OACC_KERNELS_LOOP,
	EXEC_OACC_PARALLEL, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ENTER_DATA,
	EXEC_OACC_EXIT_DATA.
	* gfortran.texi: Update for OpenACC.
	* intrinsic.texi: Likewise.
	* invoke.texi: Likewise.
	gcc/lto/
	* lto-lang.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
	New macros.
	* lto.c: Include "gomp-constants.h".
	gcc/testsuite/
	* lib/target-supports.exp (check_effective_target_fopenacc): New
	procedure.
	* g++.dg/goacc-gomp/goacc-gomp.exp: New file.
	* g++.dg/goacc/goacc.exp: Likewise.
	* gcc.dg/goacc-gomp/goacc-gomp.exp: Likewise.
	* gcc.dg/goacc/goacc.exp: Likewise.
	* gfortran.dg/goacc/goacc.exp: Likewise.
	* c-c++-common/cpp/openacc-define-1.c: New file.
	* c-c++-common/cpp/openacc-define-2.c: Likewise.
	* c-c++-common/cpp/openacc-define-3.c: Likewise.
	* c-c++-common/goacc-gomp/nesting-1.c: Likewise.
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/acc_on_device-2-off.c: Likewise.
	* c-c++-common/goacc/acc_on_device-2.c: Likewise.
	* c-c++-common/goacc/asyncwait-1.c: Likewise.
	* c-c++-common/goacc/cache-1.c: Likewise.
	* c-c++-common/goacc/clauses-fail.c: Likewise.
	* c-c++-common/goacc/collapse-1.c: Likewise.
	* c-c++-common/goacc/data-1.c: Likewise.
	* c-c++-common/goacc/data-2.c: Likewise.
	* c-c++-common/goacc/data-clause-duplicate-1.c: Likewise.
	* c-c++-common/goacc/deviceptr-1.c: Likewise.
	* c-c++-common/goacc/deviceptr-2.c: Likewise.
	* c-c++-common/goacc/deviceptr-3.c: Likewise.
	* c-c++-common/goacc/if-clause-1.c: Likewise.
	* c-c++-common/goacc/if-clause-2.c: Likewise.
	* c-c++-common/goacc/kernels-1.c: Likewise.
	* c-c++-common/goacc/loop-1.c: Likewise.
	* c-c++-common/goacc/loop-private-1.c: Likewise.
	* c-c++-common/goacc/nesting-1.c: Likewise.
	* c-c++-common/goacc/nesting-data-1.c: Likewise.
	* c-c++-common/goacc/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/parallel-1.c: Likewise.
	* c-c++-common/goacc/pcopy.c: Likewise.
	* c-c++-common/goacc/pcopyin.c: Likewise.
	* c-c++-common/goacc/pcopyout.c: Likewise.
	* c-c++-common/goacc/pcreate.c: Likewise.
	* c-c++-common/goacc/pragma_context.c: Likewise.
	* c-c++-common/goacc/present-1.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.
	* c-c++-common/goacc/sb-1.c: Likewise.
	* c-c++-common/goacc/sb-2.c: Likewise.
	* c-c++-common/goacc/sb-3.c: Likewise.
	* c-c++-common/goacc/update-1.c: Likewise.
	* gcc.dg/goacc/acc_on_device-1.c: Likewise.
	* gfortran.dg/goacc/acc_on_device-1.f95: Likewise.
	* gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise.
	* gfortran.dg/goacc/acc_on_device-2.f95: Likewise.
	* gfortran.dg/goacc/assumed.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-1.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-2.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-3.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-4.f95: Likewise.
	* gfortran.dg/goacc/branch.f95: Likewise.
	* gfortran.dg/goacc/cache-1.f95: Likewise.
	* gfortran.dg/goacc/coarray.f95: Likewise.
	* gfortran.dg/goacc/continuation-free-form.f95: Likewise.
	* gfortran.dg/goacc/cray.f95: Likewise.
	* gfortran.dg/goacc/critical.f95: Likewise.
	* gfortran.dg/goacc/data-clauses.f95: Likewise.
	* gfortran.dg/goacc/data-tree.f95: Likewise.
	* gfortran.dg/goacc/declare-1.f95: Likewise.
	* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
	* gfortran.dg/goacc/fixed-1.f: Likewise.
	* gfortran.dg/goacc/fixed-2.f: Likewise.
	* gfortran.dg/goacc/fixed-3.f: Likewise.
	* gfortran.dg/goacc/fixed-4.f: Likewise.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.
	* gfortran.dg/goacc/if.f95: Likewise.
	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
	* gfortran.dg/goacc/list.f95: Likewise.
	* gfortran.dg/goacc/literal.f95: Likewise.
	* gfortran.dg/goacc/loop-1.f95: Likewise.
	* gfortran.dg/goacc/loop-2.f95: Likewise.
	* gfortran.dg/goacc/loop-3.f95: Likewise.
	* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
	* gfortran.dg/goacc/omp.f95: Likewise.
	* gfortran.dg/goacc/parallel-kernels-clauses.f95: Likewise.
	* gfortran.dg/goacc/parallel-kernels-regions.f95: Likewise.
	* gfortran.dg/goacc/parallel-tree.f95: Likewise.
	* gfortran.dg/goacc/parameter.f95: Likewise.
	* gfortran.dg/goacc/private-1.f95: Likewise.
	* gfortran.dg/goacc/private-2.f95: Likewise.
	* gfortran.dg/goacc/private-3.f95: Likewise.
	* gfortran.dg/goacc/pure-elemental-procedures.f95: Likewise.
	* gfortran.dg/goacc/reduction-2.f95: Likewise.
	* gfortran.dg/goacc/reduction.f95: Likewise.
	* gfortran.dg/goacc/routine-1.f90: Likewise.
	* gfortran.dg/goacc/routine-2.f90: Likewise.
	* gfortran.dg/goacc/sentinel-free-form.f95: Likewise.
	* gfortran.dg/goacc/several-directives.f95: Likewise.
	* gfortran.dg/goacc/sie.f95: Likewise.
	* gfortran.dg/goacc/subarrays.f95: Likewise.
	* gfortran.dg/gomp/map-1.f90: Likewise.
	* gfortran.dg/openacc-define-1.f90: Likewise.
	* gfortran.dg/openacc-define-2.f90: Likewise.
	* gfortran.dg/openacc-define-3.f90: Likewise.
	* g++.dg/gomp/block-1.C: Update for changed compiler output.
	* g++.dg/gomp/block-2.C: Likewise.
	* g++.dg/gomp/block-3.C: Likewise.
	* g++.dg/gomp/block-5.C: Likewise.
	* g++.dg/gomp/target-1.C: Likewise.
	* g++.dg/gomp/target-2.C: Likewise.
	* g++.dg/gomp/taskgroup-1.C: Likewise.
	* g++.dg/gomp/teams-1.C: Likewise.
	* gcc.dg/cilk-plus/jump-openmp.c: Likewise.
	* gcc.dg/cilk-plus/jump.c: Likewise.
	* gcc.dg/gomp/block-1.c: Likewise.
	* gcc.dg/gomp/block-10.c: Likewise.
	* gcc.dg/gomp/block-2.c: Likewise.
	* gcc.dg/gomp/block-3.c: Likewise.
	* gcc.dg/gomp/block-4.c: Likewise.
	* gcc.dg/gomp/block-5.c: Likewise.
	* gcc.dg/gomp/block-6.c: Likewise.
	* gcc.dg/gomp/block-7.c: Likewise.
	* gcc.dg/gomp/block-8.c: Likewise.
	* gcc.dg/gomp/block-9.c: Likewise.
	* gcc.dg/gomp/target-1.c: Likewise.
	* gcc.dg/gomp/target-2.c: Likewise.
	* gcc.dg/gomp/taskgroup-1.c: Likewise.
	* gcc.dg/gomp/teams-1.c: Likewise.
	include/
	* gomp-constants.h: New file.
	libgomp/
	* Makefile.am (search_path): Add $(top_srcdir)/../include.
	(libgomp_la_SOURCES): Add splay-tree.c, libgomp-plugin.c,
	oacc-parallel.c, oacc-host.c, oacc-init.c, oacc-mem.c,
	oacc-async.c, oacc-plugin.c, oacc-cuda.c.
	[USE_FORTRAN] (libgomp_la_SOURCES): Add openacc.f90.
	Include $(top_srcdir)/plugin/Makefrag.am.
	(nodist_libsubinclude_HEADERS): Add openacc.h.
	[USE_FORTRAN] (nodist_finclude_HEADERS): Add openacc_lib.h,
	openacc.f90, openacc.mod, openacc_kinds.mod.
	(omp_lib.mod): Generalize into...
	(%.mod): ... this new rule.
	(openacc_kinds.mod, openacc.mod): New rules.
	* plugin/configfrag.ac: New file.
	* configure.ac: Move plugin/offloading support into it.  Include
	it.  Instantiate testsuite/libgomp-test-support.pt.exp.
	* plugin/Makefrag.am: New file.
	* testsuite/Makefile.am (OFFLOAD_TARGETS)
	(OFFLOAD_ADDITIONAL_OPTIONS, OFFLOAD_ADDITIONAL_LIB_PATHS): Don't
	export.
	(libgomp-test-support.exp): New rule.
	(all-local): Depend on it.
	* Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.
	* config.h.in: Likewise.
	* configure: Likewise.
	* configure.tgt: Harden shell syntax.
	* env.c: Include "oacc-int.h".
	(parse_acc_device_type): New function.
	(gomp_debug_var, goacc_device_type, goacc_device_num): New
	variables.
	(initialize_env): Initialize those.  Call
	goacc_runtime_initialize.
	* error.c (gomp_vdebug, gomp_debug, gomp_vfatal): New functions.
	(gomp_fatal): Call gomp_vfatal.
	* libgomp.h: Include "libgomp-plugin.h" and <stdarg.h>.
	(gomp_debug_var, goacc_device_type, goacc_device_num, gomp_vdebug)
	(gomp_debug, gomp_verror, gomp_vfatal, gomp_init_targets_once)
	(splay_tree_node, splay_tree, splay_tree_key)
	(struct target_mem_desc, struct splay_tree_key_s)
	(struct gomp_memory_mapping, struct acc_dispatch_t)
	(struct gomp_device_descr, gomp_acc_insert_pointer)
	(gomp_acc_remove_pointer, target_mem_desc, gomp_copy_from_async)
	(gomp_unmap_vars, gomp_init_device, gomp_init_tables)
	(gomp_free_memmap, gomp_fini_device): New declarations.
	(gomp_vdebug, gomp_debug): New macros.
	Include "splay-tree.h".
	* libgomp.map (OACC_2.0): New symbol version.  Use for
	acc_get_num_devices, acc_get_num_devices_h_, acc_set_device_type,
	acc_set_device_type_h_, acc_get_device_type,
	acc_get_device_type_h_, acc_set_device_num, acc_set_device_num_h_,
	acc_get_device_num, acc_get_device_num_h_, acc_async_test,
	acc_async_test_h_, acc_async_test_all, acc_async_test_all_h_,
	acc_wait, acc_wait_h_, acc_wait_async, acc_wait_async_h_,
	acc_wait_all, acc_wait_all_h_, acc_wait_all_async,
	acc_wait_all_async_h_, acc_init, acc_init_h_, acc_shutdown,
	acc_shutdown_h_, acc_on_device, acc_on_device_h_, acc_malloc,
	acc_free, acc_copyin, acc_copyin_32_h_, acc_copyin_64_h_,
	acc_copyin_array_h_, acc_present_or_copyin,
	acc_present_or_copyin_32_h_, acc_present_or_copyin_64_h_,
	acc_present_or_copyin_array_h_, acc_create, acc_create_32_h_,
	acc_create_64_h_, acc_create_array_h_, acc_present_or_create,
	acc_present_or_create_32_h_, acc_present_or_create_64_h_,
	acc_present_or_create_array_h_, acc_copyout, acc_copyout_32_h_,
	acc_copyout_64_h_, acc_copyout_array_h_, acc_delete,
	acc_delete_32_h_, acc_delete_64_h_, acc_delete_array_h_,
	acc_update_device, acc_update_device_32_h_,
	acc_update_device_64_h_, acc_update_device_array_h_,
	acc_update_self, acc_update_self_32_h_, acc_update_self_64_h_,
	acc_update_self_array_h_, acc_map_data, acc_unmap_data,
	acc_deviceptr, acc_hostptr, acc_is_present, acc_is_present_32_h_,
	acc_is_present_64_h_, acc_is_present_array_h_,
	acc_memcpy_to_device, acc_memcpy_from_device,
	acc_get_current_cuda_device, acc_get_current_cuda_context,
	acc_get_cuda_stream, acc_set_cuda_stream.
	(GOACC_2.0): New symbol version.  Use for GOACC_data_end,
	GOACC_data_start, GOACC_enter_exit_data, GOACC_parallel,
	GOACC_update, GOACC_wait, GOACC_get_thread_num,
	GOACC_get_num_threads.
	(GOMP_PLUGIN_1.0): New symbol version.  Use for
	GOMP_PLUGIN_malloc, GOMP_PLUGIN_malloc_cleared,
	GOMP_PLUGIN_realloc, GOMP_PLUGIN_debug, GOMP_PLUGIN_error,
	GOMP_PLUGIN_fatal, GOMP_PLUGIN_async_unmap_vars,
	GOMP_PLUGIN_acc_thread.
	* libgomp.texi: Update for OpenACC changes, and GOMP_DEBUG
	environment variable.
	* libgomp_g.h (GOACC_data_start, GOACC_data_end)
	(GOACC_enter_exit_data, GOACC_parallel, GOACC_update, GOACC_wait)
	(GOACC_get_num_threads, GOACC_get_thread_num): New declarations.
	* splay-tree.h (splay_tree_lookup, splay_tree_insert)
	(splay_tree_remove): New declarations.
	(rotate_left, rotate_right, splay_tree_splay, splay_tree_insert)
	(splay_tree_remove, splay_tree_lookup): Move into...
	* splay-tree.c: ... this new file.
	* target.c: Include "oacc-plugin.h", "oacc-int.h", <assert.h>.
	(splay_tree_node, splay_tree, splay_tree_key)
	(struct target_mem_desc, struct splay_tree_key_s)
	(struct gomp_device_descr): Don't declare.
	(num_devices_openmp): New variable.
	(gomp_get_num_devices ): Use it.
	(gomp_init_targets_once): New function.
	(gomp_get_num_devices ): Use it.
	(get_kind, gomp_copy_from_async, gomp_free_memmap)
	(gomp_fini_device, gomp_register_image_for_device): New functions.
	(gomp_map_vars): Add devaddrs parameter.
	(gomp_update): Add mm parameter.
	(gomp_init_device): Move most of it into...
	(gomp_init_tables): ... this new function.
	(gomp_register_images_for_device): Remove function.
	(splay_compare, gomp_map_vars, gomp_unmap_vars, gomp_init_device):
	Make them hidden instead of static.
	(gomp_map_vars_existing, gomp_map_vars, gomp_unmap_vars)
	(gomp_update, gomp_init_device, GOMP_target, GOMP_target_data)
	(GOMP_target_end_data, GOMP_target_update)
	(gomp_load_plugin_for_device, gomp_target_init): Update for
	OpenACC changes.
	* oacc-async.c: New file.
	* oacc-cuda.c: Likewise.
	* oacc-host.c: Likewise.
	* oacc-init.c: Likewise.
	* oacc-int.h: Likewise.
	* oacc-mem.c: Likewise.
	* oacc-parallel.c: Likewise.
	* oacc-plugin.c: Likewise.
	* oacc-plugin.h: Likewise.
	* oacc-ptx.h: Likewise.
	* openacc.f90: Likewise.
	* openacc.h: Likewise.
	* openacc_lib.h: Likewise.
	* plugin/plugin-host.c: Likewise.
	* plugin/plugin-nvptx.c: Likewise.
	* libgomp-plugin.c: Likewise.
	* libgomp-plugin.h: Likewise.
	* libgomp_target.h: Remove file after merging content into the
	former file.  Update all users.
	* testsuite/lib/libgomp.exp: Load libgomp-test-support.exp.
	(offload_targets_s, offload_targets_s_openacc): New variables.
	(check_effective_target_openacc_nvidia_accel_present)
	(check_effective_target_openacc_nvidia_accel_selected): New
	procedures.
	(libgomp_init): Update for OpenACC changes.
	* testsuite/libgomp-test-support.exp.in: New file.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/cache-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-1.c: Likewise.
	* 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/data-already-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-10.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-11.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-12.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-19.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-26.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-27.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-31.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-33.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-35.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-36.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-37.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-39.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-40.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-41.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-45.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-46.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-49.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-50.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-51.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-55.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-56.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-57.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-58.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-59.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-6.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-60.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-61.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-62.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-63.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-64.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-65.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-66.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-67.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-68.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-7.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-c-c++-common/lib-84.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-86.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-87.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-88.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-89.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-9.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-90.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-92.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/offset-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-empty.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pointer-align-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c:
	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: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Likewise.
	* testsuite/libgomp.oacc-fortran/abort-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/abort-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/asyncwait-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-5.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-6.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-7.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/collapse-8.f90: 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-fortran/data-already-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-10.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-5.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-6.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-7.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-8.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/map-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/openacc_version-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/pointer-align-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/pset-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/subarrays-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/subarrays-2.f90: Likewise.
	liboffloadmic/
	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_name)
	(GOMP_OFFLOAD_get_caps, GOMP_OFFLOAD_fini_device): New functions.

Co-Authored-By: Bernd Schmidt <bernds@codesourcery.com>
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Dmitry Bocharnikov <dmitry.b@samsung.com>
Co-Authored-By: Evgeny Gavrin <e.gavrin@samsung.com>
Co-Authored-By: Ilmir Usmanov <i.usmanov@samsung.com>
Co-Authored-By: Jakub Jelinek <jakub@redhat.com>
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Tobias Burnus <burnus@net-b.de>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r219682
2015-01-15 21:11:12 +01:00
Thomas Schwinge f9a0eca40d libgomp: Fix "intelmic" offloading in build-tree testing.
libgomp/
	* testsuite/lib/libgomp.exp (libgomp_init): Correctly match
	"intelmic" in $offload_targets.

From-SVN: r219348
2015-01-08 17:01:24 +01:00
Jakub Jelinek 5624e564d2 Update copyright years.
From-SVN: r219188
2015-01-05 13:33:28 +01:00
Kyrylo Tkachov 3b41b58357 [libgomp][testsuite] Explicitly include target-utils.exp (fix libgomp testsuite ERROR)
* testsuite/lib/libgomp.exp: Load target-utils.exp.
	Move load of target-supportes.exp earlier.

From-SVN: r218662
2014-12-12 10:31:44 +00:00
Ilya Verbin 2354caecd9 Force output of vars with "omp declare target" attribute in gcc/varpool.c
gcc/
	* varpool.c (varpool_node::get_create): Force output of vars with
	"omp declare target" attribute.
libgomp/
	* testsuite/libgomp.c/target-9.c: New test.

From-SVN: r218607
2014-12-10 20:52:10 +00:00
Andrey Turetskiy 476ff78736 omp-low.c (lower_omp_critical): Mark critical sections inside target functions as offloadable.
gcc/
	* omp-low.c (lower_omp_critical): Mark critical sections
	inside target functions as offloadable.
libgomp/
	* testsuite/libgomp.c/target-critical-1.c: New test.

Co-Authored-By: Ilya Verbin <ilya.verbin@intel.com>

From-SVN: r218158
2014-11-28 13:59:49 +00:00
Jakub Jelinek 31a30e6234 e.53.4.c: Add -DITESTITERS=20 to dg-options unless expensive testing is on.
* testsuite/libgomp.c/examples-4/e.53.4.c: Add -DITESTITERS=20
	to dg-options unless expensive testing is on.  
	(TESTITERS): Define to N if not defined.
	(main): Use TESTITERS instead of N.
	* testsuite/libgomp.c/examples-4/e.55.1.c: Define CHUNKSZ from
	dg-additional-options depending on whether expensive testing is on.
	* testsuite/libgomp.fortran/examples-4/e.55.1.f90 (e_55_1_mod):
	Decrease N to 100000 and CHUNKSZ to 10000.

From-SVN: r218095
2014-11-26 20:36:36 +01:00
Jakub Jelinek 3f9e8f13a3 re PR fortran/63938 (OpenMP atomic update does not protect access to automatic array)
PR fortran/63938
	* trans-openmp.c (gfc_trans_omp_atomic): Make sure lhsaddr is
	simple enough for goa_lhs_expr_p.

	* libgomp.fortran/pr63938-1.f90: New test.
	* libgomp.fortran/pr63938-2.f90: New test.

From-SVN: r218031
2014-11-25 00:08:26 +01:00
Uros Bizjak 4a19f719d6 e.53.5.c: Require vect_simd_clones effective target.
* testsuite/libgomp.c/examples-4/e.53.5.c: Require
	vect_simd_clones effective target.
	* testsuite/libgomp.fortran/examples-4/e.53.5.f90: Ditto.

From-SVN: r217760
2014-11-19 11:30:29 +01:00
Jakub Jelinek b1bf6c522f e.54.2.c (main): Use N / 8 instead of 32 as block_size.
* libgomp.c/examples-4/e.54.2.c (main): Use N / 8 instead
	of 32 as block_size.
	* libgomp.fortran/examples-4/e.54.2.f90 (e_54_1): Use n / 8
	instead of 32 as block_size.

From-SVN: r217579
2014-11-14 18:19:55 +01:00
Andrey Turetskiy d64ae6140f [PATCH 4/4] OpenMP 4.0 offloading to Intel MIC: non-fallback testing.
libgomp/
	* Makefile.in: Regenerate.
	* configure: Regenerate.
	* configure.ac: Set up offload_additional_options,
	offload_additional_lib_paths and offload_targets.
	* testsuite/Makefile.am: Export environment variables: OFFLOAD_TARGETS,
	OFFLOAD_ADDITIONAL_OPTIONS, OFFLOAD_ADDITIONAL_LIB_PATHS.
	* testsuite/Makefile.in: Regenerate.
	* testsuite/lib/libgomp.exp (libgomp_init): Append
	offload_additional_lib_paths to LD_LIBRARY_PATH.  Append
	offload_additional_options to ALWAYS_CFLAGS.  Append liboffloadmic
	build directory to LD_LIBRARY_PATH for intelmic offload targets.

Co-Authored-By: Ilya Verbin <ilya.verbin@intel.com>

From-SVN: r217500
2014-11-13 14:07:09 +00:00
Andrey Turetskiy 122d7303a0 [PATCH 7/7] OpenMP 4.0 offloading infrastructure: testsuite.
libgomp/
	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_device): New.
	* testsuite/libgomp.c++/c++.exp: Include tests from subdirectories.
	* testsuite/libgomp.c++/examples-4/e.51.5.C: New test.
	* testsuite/libgomp.c++/examples-4/e.53.2.C: Ditto.
	* testsuite/libgomp.c/examples-4/e.50.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.50.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.50.3.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.50.4.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.50.5.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.3.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.4.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.6.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.51.7.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.52.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.52.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.53.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.53.3.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.53.4.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.53.5.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.54.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.54.3.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.54.4.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.54.5.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.54.6.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.55.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.55.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.56.3.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.56.4.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.57.1.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.57.2.c: Ditto.
	* testsuite/libgomp.c/examples-4/e.57.3.c: Ditto.
	* testsuite/libgomp.c/target-7.c: Fix test.
	* testsuite/libgomp.fortran/examples-4/e.50.1.f90: New test.
	* testsuite/libgomp.fortran/examples-4/e.50.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.50.3.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.50.4.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.50.5.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.1.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.3.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.4.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.5.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.6.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.51.7.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.52.1.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.52.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.53.1.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.53.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.53.3.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.53.4.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.53.5.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.54.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.54.3.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.54.4.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.54.5.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.54.6.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.55.1.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.55.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.56.3.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.56.4.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.57.1.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.57.2.f90: Ditto.
	* testsuite/libgomp.fortran/examples-4/e.57.3.f90: Ditto.

Co-Authored-By: Ilya Tocar <ilya.tocar@intel.com>
Co-Authored-By: Ilya Verbin <ilya.verbin@intel.com>
Co-Authored-By: Kirill Yukhin <kirill.yukhin@intel.com>

From-SVN: r217494
2014-11-13 13:56:22 +00:00
Marek Polacek c1d62412c3 affinity-1.c: Include <sys/wait.h>.
* testsuite/libgomp.c/affinity-1.c: Include <sys/wait.h>.
	* testsuite/libgomp.c/nqueens-1.c: Include <stdlib.h>.
	* testsuite/libgomp.c/thread-limit-1.c: Include <omp.h>
	* testsuite/libgomp.c/thread-limit-2.c: Likewise.

From-SVN: r215927
2014-10-06 11:54:24 +00:00
Marek Polacek 44a0c6cbeb affinity-1.c: Fix implicit declarations.
* testsuite/libgomp.c/affinity-1.c: Fix implicit declarations.
	* testsuite/libgomp.c/nqueens-1.c: Likewise.
	* testsuite/libgomp.c/pr26943-3.c: Likewise.
	* testsuite/libgomp.c/pr26943-4.c: Likewise.
	* testsuite/libgomp.c/pr36802-2.c: Likewise.
	* testsuite/libgomp.c/pr36802-3.c: Likewise.
	* testsuite/libgomp.c/thread-limit-1.c: Likewise.
	* testsuite/libgomp.c/thread-limit-2.c: Likewise.
	* testsuite/libgomp.c/appendix-a/a.15.1.c: Include <omp.h>.
	* testsuite/libgomp.c/omp-loop02.c: Fix defaulting to int.
	* testsuite/libgomp.c/omp-parallel-for.c: Likewise.
	* testsuite/libgomp.c/omp-parallel-if.c: Likewise.
	* testsuite/libgomp.c/omp-single-1.c: Likewise.
	* testsuite/libgomp.c/omp-single-2.c: Likewise.
	* testsuite/libgomp.c/omp_matvec.c: Likewise.
	* testsuite/libgomp.c/omp_workshare3.c: Likewise.
	* testsuite/libgomp.c/omp_workshare4.c: Likewise.
	* testsuite/libgomp.c/shared-1.c: Fix defaulting to int.  Fix implicit
	declarations.

From-SVN: r215922
2014-10-06 10:20:45 +00:00
Jakub Jelinek 5771c3915b re PR libgomp/61200 (internal compiler error: Segmentation fault, assert & openmp)
PR libgomp/61200
	* omp-low.c (taskreg_contexts): New variable.
	(scan_omp_parallel): Push newly created context into taskreg_contexts
	vector and move record layout code to finish_taskreg_scan.
	(scan_omp_task): Likewise.
	(finish_taskreg_scan): New function.
	(execute_lower_omp): Call finish_taskreg_scan on all taskreg_contexts
	vector elements and release it.

	* c-c++-common/gomp/pr61200.c: New test.

	* testsuite/libgomp.c/pr61200.c: New test.

From-SVN: r215835
2014-10-03 09:29:42 +02:00
Jakub Jelinek bce16b887f re PR c++/63248 (Crash when OpenMP target's array section handling is used with templates)
PR c++/63248
	* semantics.c (finish_omp_clauses): Don't call cp_omp_mappable_type
	on type of type dependent expressions, and don't call it if
	handle_omp_array_sections has kept TREE_LIST because something
	was type dependent.
	* pt.c (tsubst_expr) <case OMP_TARGET, case OMP_TARGET_DATA>:
	Use keep_next_level, begin_omp_structured_block and
	finish_omp_structured_block instead of push_stmt_list and
	pop_stmt_list.
libgomp/
	* testsuite/libgomp.c++/pr63248.C: New test.

From-SVN: r215359
2014-09-18 18:43:28 +02:00
Jakub Jelinek 3696163cb4 task.c (GOMP_taskgroup_end): If taskgroup->num_children is not zero...
* task.c (GOMP_taskgroup_end): If taskgroup->num_children
	is not zero, but taskgroup->children is NULL and there are
	any task->children, schedule those instead of waiting.
	* testsuite/libgomp.c/depend-6.c: New test.
	* testsuite/libgomp.c/depend-7.c: New test.
	* testsuite/libgomp.c/depend-8.c: New test.
	* testsuite/libgomp.c/depend-9.c: New test.
	* testsuite/libgomp.c/depend-10.c: New test.

From-SVN: r213592
2014-08-04 17:45:50 +02:00
Jakub Jelinek 0494285ab0 libgomp.h (struct gomp_task_depend_entry): Add redundant_out field.
* libgomp.h (struct gomp_task_depend_entry): Add redundant_out field.
	(struct gomp_taskwait): New type.
	(struct gomp_task): Add taskwait and parent_depends_on, remove
	in_taskwait and taskwait_sem fields.
	(gomp_finish_task): Don't destroy taskwait_sem.
	* task.c (gomp_init_task): Don't init in_taskwait, instead init
	taskwait and parent_depends_on.
	(GOMP_task): For if (0) tasks with depend clause that depend on
	earlier tasks don't defer them, instead call
	gomp_task_maybe_wait_for_dependencies to wait for the dependencies.
	Initialize redundant_out field, for redundant out entries just
	move them at the end of linked list instead of removing them
	completely, and set redundant_out flag instead of redundant.
	(gomp_task_run_pre): Update last_parent_depends_on if scheduling
	that task.
	(gomp_task_run_post_handle_dependers): If parent is in
	gomp_task_maybe_wait_for_dependencies and newly runnable task
	is not parent_depends_on, queue it in parent->children linked
	list after all runnable tasks with parent_depends_on set.
	Adjust for addition of taskwait indirection.
	(gomp_task_run_post_remove_parent): If parent is in
	gomp_task_maybe_wait_for_dependencies and task to be removed
	is parent_depends_on, decrement n_depend and if needed awake
	parent.  Adjust for addition of taskwait indirection.
	(GOMP_taskwait): Adjust for addition of taskwait indirection.
	(gomp_task_maybe_wait_for_dependencies): New function.
	* testsuite/libgomp.c/depend-5.c: New test.

From-SVN: r213408
2014-08-01 10:05:13 +02:00
Tobias Burnus 60ab4b449c pr34020.f90: Make compile with TS 18508/Fortran 2015
2014-07-13  Tobias Burnus  <burnus@net-b.de>

        * testsuite/libgomp.fortran/pr34020.f90: Make compile
        with TS 18508/Fortran 2015

From-SVN: r212490
2014-07-13 08:40:16 +02:00
Marek Polacek 773ec47fe6 re PR c/6940 (taking sizeof array parameter should trigger a warning)
PR c/6940
	* doc/invoke.texi: Document -Wsizeof-array-argument.
c-family/
	* c.opt (Wsizeof-array-argument): New option.
c/
	* c-decl.c (grokdeclarator): Set C_ARRAY_PARAMETER.
	* c-tree.h (C_ARRAY_PARAMETER): Define.
	* c-typeck.c (c_expr_sizeof_expr): Warn when using sizeof on an array
	function parameter.
cp/
	* cp-tree.h (DECL_ARRAY_PARAMETER_P): Define.
	* decl.c (grokdeclarator): Set DECL_ARRAY_PARAMETER_P.
	* typeck.c (cxx_sizeof_expr): Warn when using sizeof on an array
	function parameter.
testsuite/
	* c-c++-common/Wsizeof-pointer-memaccess1.c: Use
	-Wno-sizeof-array-argument.
	* c-c++-common/Wsizeof-pointer-memaccess2.c: Likewise.
	* g++.dg/warn/Wsizeof-pointer-memaccess-1.C: Likewise.
	* gcc.dg/Wsizeof-pointer-memaccess1.c: Likewise.
	* g++.dg/torture/Wsizeof-pointer-memaccess1.C: Likewise.
	* g++.dg/torture/Wsizeof-pointer-memaccess2.C: Likewise.
	* gcc.dg/torture/Wsizeof-pointer-memaccess1.c: Likewise.
	* c-c++-common/sizeof-array-argument.c: New test.
	* gcc.dg/vla-5.c: Add dg-warnings.
../libgomp/	* testsuite/libgomp.c/appendix-a/a.29.1.c (f): Add dg-warnings.

From-SVN: r212312
2014-07-06 19:00:10 +00:00
Thomas Schwinge a03bd0056c Normalize interface for all *-dg-runtest.
gcc/testsuite/
	* lib/g++-dg.exp (g++-dg-runtest): Change interface to match
	dg-runtest's.  Adapt all callers.
	* lib/gcc-dg.exp (gcc-dg-runtest): Likewise.
	* lib/gfortran-dg.exp (gfortran-dg-runtest): Likewise.
	* lib/go-dg.exp (go-dg-runtest): Likewise.
	* lib/obj-c++-dg.exp (obj-c++-dg-runtest): Likewise.
	* lib/objc-dg.exp (objc-dg-runtest): Likewise.
	libffi/
	* testsuite/lib/libffi.exp (libffi-dg-runtest): Change interface
	match to dg-runtest's.

From-SVN: r212278
2014-07-04 07:28:17 +02:00
Jakub Jelinek f707da16f7 libgomp.exp (libgomp_target_compile): If $source matches regex $lang_source_re, add $lang_include_flags to options.
* testsuite/lib/libgomp.exp (libgomp_target_compile): If $source
	matches regex $lang_source_re, add $lang_include_flags to options.
	* testsuite/libgomp.c/c.exp: Unset lang_include_flags.
	* testsuite/libgomp.c++/c++.exp: Likewise.
	* testsuite/libgomp.fortran/fortran.exp: Likewise.  Set lang_source_re
	and lang_include_flags instead of adding -fintrinsic-modules-path= to
	ALWAYS_CFLAGS.
	* testsuite/libgomp.graphite/graphite.exp: Unset lang_include_flags.

From-SVN: r212268
2014-07-03 20:45:35 +02:00
Thomas Schwinge 07b6c04483 Explain gfortran-dg-runtest usage in libgomp Fortran testing.
libgomp/
	* testsuite/libgomp.fortran/fortran.exp: Explain
	gfortran-dg-runtest usage.

From-SVN: r212266
2014-07-03 17:23:43 +02:00
Jakub Jelinek da6f124d8a langhooks-def.h (LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR): Define.
* langhooks-def.h (LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR): Define.
	(LANG_HOOKS_DECLS): Add it.
	* gimplify.c (gimplify_omp_for): Make sure OMP_CLAUSE_LINEAR_STEP
	has correct type.
	* tree.h (OMP_CLAUSE_LINEAR_ARRAY): Define.
	* langhooks.h (struct lang_hooks_for_decls): Add
	omp_clause_linear_ctor hook.
	* omp-low.c (lower_rec_input_clauses): Set max_vf even if
	OMP_CLAUSE_LINEAR_ARRAY is set.  Don't fold_convert
	OMP_CLAUSE_LINEAR_STEP.  For OMP_CLAUSE_LINEAR_ARRAY in
	combined simd loop use omp_clause_linear_ctor hook.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): Make sure
	OMP_CLAUSE_LINEAR_STEP has correct type.
gcc/cp/
	* semantics.c (finish_omp_clauses): Make sure
	OMP_CLAUSE_LINEAR_STEP has correct type.
gcc/fortran/
	* trans.h (gfc_omp_clause_linear_ctor): New prototype.
	* trans-openmp.c (gfc_omp_linear_clause_add_loop,
	gfc_omp_clause_linear_ctor): New functions.
	(gfc_trans_omp_clauses): Make sure OMP_CLAUSE_LINEAR_STEP has
	correct type.  Set OMP_CLAUSE_LINEAR_ARRAY flag if needed.
	* f95-lang.c (LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR): Redefine.
libgomp/
	* testsuite/libgomp.fortran/simd5.f90: New test.
	* testsuite/libgomp.fortran/simd6.f90: New test.
	* testsuite/libgomp.fortran/simd7.f90: New test.

From-SVN: r211971
2014-06-25 11:16:12 +02:00
Jakub Jelinek 56ad0e3820 gimplify.c (gimplify_omp_for): For #pragma omp for simd iterator not mentioned in clauses use private clause if...
* gimplify.c (gimplify_omp_for): For #pragma omp for simd iterator
	not mentioned in clauses use private clause if the iterator is
	declared in #pragma omp for simd, and when adding lastprivate
	instead, add it to the outer #pragma omp for too.  Diagnose
	if the variable is private in outer context.  For simd collapse > 1
	loops, replace all iterators with temporaries.
	* omp-low.c (lower_rec_input_clauses): Handle LINEAR clause the
	same even in collapse > 1 loops.
gcc/c/
	* c-parser.c (c_parser_omp_for_loop): For
	#pragma omp parallel for simd move lastprivate clause from parallel
	to for rather than simd.
gcc/cp/
	* parser.c (cp_parser_omp_for_loop): For
	#pragma omp parallel for simd move lastprivate clause from parallel
	to for rather than simd.
libgomp/
	* testsuite/libgomp.c/for-2.c: Define SC to static for
	#pragma omp for simd testing.
	* testsuite/libgomp.c/for-2.h (SC): Define if not defined.
	(N(f5), N(f6), N(f7), N(f8), N(f10), N(f12), N(f14)): Use
	SC macro.
	* testsuite/libgomp.c/simd-14.c: New test.
	* testsuite/libgomp.c/simd-15.c: New test.
	* testsuite/libgomp.c/simd-16.c: New test.
	* testsuite/libgomp.c/simd-17.c: New test.
	* testsuite/libgomp.c++/for-10.C: Define SC to static for
	#pragma omp for simd testing.
	* testsuite/libgomp.c++/simd10.C: New test.
	* testsuite/libgomp.c++/simd11.C: New test.
	* testsuite/libgomp.c++/simd12.C: New test.
	* testsuite/libgomp.c++/simd13.C: New test.

From-SVN: r211930
2014-06-24 09:53:52 +02:00
Jakub Jelinek b46ebd6c7b gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP, [...]): Make sure OMP_CLAUSE_SIZE is non-NULL.
* gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP,
	OMP_CLAUSE_TO, OMP_CLAUSE_FROM): Make sure OMP_CLAUSE_SIZE is
	non-NULL.
	<case OMP_CLAUSE_ALIGNED>: Gimplify OMP_CLAUSE_ALIGNED_ALIGNMENT.
	(gimplify_adjust_omp_clauses_1): Make sure OMP_CLAUSE_SIZE is
	non-NULL.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-low.c (lower_rec_simd_input_clauses,
	lower_rec_input_clauses, expand_omp_simd): Handle non-constant
	safelen the same as safelen(1).
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Handle OMP_CLAUSE_ALIGNED.  For
	OMP_CLAUSE_{MAP,TO,FROM} if not decl use walk_tree.
	(convert_nonlocal_reference_stmt, convert_local_reference_stmt):
	Fixup handling of GIMPLE_OMP_TARGET.
	(convert_tramp_reference_stmt, convert_gimple_call): Handle
	GIMPLE_OMP_TARGET.
gcc/fortran/
	* dump-parse-tree.c (show_omp_namelist): Use n->udr->udr instead
	of n->udr.
	* f95-lang.c (gfc_init_builtin_functions): Initialize
	BUILT_IN_ASSUME_ALIGNED.
	* gfortran.h (gfc_omp_namelist): Change udr field type to
	struct gfc_omp_namelist_udr.
	(gfc_omp_namelist_udr): New type.
	(gfc_get_omp_namelist_udr): Define.
	(gfc_resolve_code): New prototype.
	* match.c (gfc_free_omp_namelist): Free name->udr.
	* module.c (intrinsics): Add INTRINSIC_USER.
	(fix_mio_expr): Likewise.
	(mio_expr): Handle INSTRINSIC_USER and non-resolved EXPR_FUNCTION.
	* openmp.c (gfc_match_omp_clauses): Adjust initialization of n->udr.
	(gfc_match_omp_declare_reduction): Treat len=: the same as len=*.
	Set attr.flavor on omp_{out,in,priv,orig} artificial variables.
	(struct resolve_omp_udr_callback_data): New type.
	(resolve_omp_udr_callback, resolve_omp_udr_callback2,
	resolve_omp_udr_clause): New functions.
	(resolve_omp_clauses): Adjust for n->udr changes, resolve UDR clauses
	here.
	(omp_udr_callback): Don't check for implicitly declared functions
	here.
	(gfc_resolve_omp_udr): Don't call gfc_resolve.  Don't check for
	implicitly declared subroutines here.
	* resolve.c (resolve_function): If value.function.isym is non-NULL,
	consider it already resolved.
	(resolve_code): Renamed to ...
	(gfc_resolve_code): ... this.  No longer static.
	(gfc_resolve_blocks, generate_component_assignments, resolve_codes):
	Adjust callers.
	* trans-openmp.c (gfc_omp_privatize_by_reference): Don't privatize
	by reference type (C_PTR) variables.
	(gfc_omp_finish_clause): Make sure OMP_CLAUSE_SIZE is non-NULL.
	(gfc_trans_omp_udr_expr): Remove.
	(gfc_trans_omp_array_reduction_or_udr): Adjust for n->udr changes.
	Don't call gfc_trans_omp_udr_expr, even for sym->attr.dimension
	expand it as assignment or subroutine call.  Don't initialize
	value.function.isym.
gcc/testsuite/
	* gfortran.dg/gomp/udr2.f90 (f7, f9): Add !$omp parallel with
	reduction clause.
	* gfortran.dg/gomp/udr4.f90 (f4): Likewise.
	Remove Label is never defined expected error.
	* gfortran.dg/gomp/udr8.f90: New test.
libgomp/
	* testsuite/libgomp.fortran/aligned1.f03: New test.
	* testsuite/libgomp.fortran/nestedfn5.f90: New test.
	* testsuite/libgomp.fortran/target7.f90: Surround loop spawning
	tasks with !$omp parallel !$omp single.
	* testsuite/libgomp.fortran/target8.f90: New test.
	* testsuite/libgomp.fortran/udr4.f90 (foo UDR, bar UDR): Adjust
	not to use trim in the combiner, instead call elemental function.
	(fn): New elemental function.
	* testsuite/libgomp.fortran/udr6.f90 (do_add, dp_add, dp_init):
	Make elemental.
	* testsuite/libgomp.fortran/udr7.f90 (omp_priv, omp_orig, omp_out,
	omp_in): Likewise.
	* testsuite/libgomp.fortran/udr12.f90: New test.
	* testsuite/libgomp.fortran/udr13.f90: New test.
	* testsuite/libgomp.fortran/udr14.f90: New test.
	* testsuite/libgomp.fortran/udr15.f90: New test.

From-SVN: r211929
2014-06-24 09:45:22 +02:00