Commit Graph

236 Commits

Author SHA1 Message Date
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
Jakub Jelinek f014c65363 gimplify.c (omp_notice_variable): If n is non-NULL and no flags change in ORT_TARGET region, don't jump to do_outer.
* gimplify.c (omp_notice_variable): If n is non-NULL
	and no flags change in ORT_TARGET region, don't jump to
	do_outer.
	(struct gimplify_adjust_omp_clauses_data): New type.
	(gimplify_adjust_omp_clauses_1): Adjust for data being
	a struct gimplify_adjust_omp_clauses_data pointer instead
	of tree *.  Pass pre_p as a new argument to
	lang_hooks.decls.omp_finish_clause hook.
	(gimplify_adjust_omp_clauses): Add pre_p argument, adjust
	splay_tree_foreach to pass both list_p and pre_p.
	(gimplify_omp_parallel, gimplify_omp_task, gimplify_omp_for,
	gimplify_omp_workshare, gimplify_omp_target_update): Adjust
	gimplify_adjust_omp_clauses callers.
	* langhooks.c (lhd_omp_finish_clause): New function.
	* langhooks-def.h (lhd_omp_finish_clause): New prototype.
	(LANG_HOOKS_OMP_FINISH_CLAUSE): Define to lhd_omp_finish_clause.
	* langhooks.h (struct lang_hooks_for_decls): Add a new
	gimple_seq * argument to omp_finish_clause hook.
	* omp-low.c (scan_sharing_clauses): Call scan_omp_op on
	non-DECL_P OMP_CLAUSE_DECL if ctx->outer.
	(scan_omp_parallel, lower_omp_for): When adding
	_LOOPTEMP_ clause var, add it to outer ctx's decl_map
	as identity.
	* tree-core.h (OMP_CLAUSE_MAP_TO_PSET): New map kind.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Handle various OpenMP 4.0 clauses.
	* tree-pretty-print.c (dump_omp_clause): Handle
	OMP_CLAUSE_MAP_TO_PSET.
gcc/cp/
	* cp-gimplify.c (cxx_omp_finish_clause): Add a gimple_seq *
	argument.
	* cp-tree.h (cxx_omp_finish_clause): Adjust prototype.
gcc/fortran/
	* cpp.c (cpp_define_builtins): Change _OPENMP macro to
	201307.
	* dump-parse-tree.c (show_omp_namelist): Add list_type
	argument.  Adjust for rop being u.reduction_op now,
	handle depend_op or map_op.
	(show_omp_node): Adjust callers.  Print some new
	OpenMP 4.0 clauses, adjust for OMP_LIST_DEPEND_{IN,OUT}
	becoming a single OMP_LIST_DEPEND.
	* f95-lang.c (gfc_handle_omp_declare_target_attribute): New
	function.
	(gfc_attribute_table): New variable.
	(LANG_HOOKS_OMP_FINISH_CLAUSE, LANG_HOOKS_ATTRIBUTE_TABLE): Redefine.
	* frontend-passes.c (gfc_code_walker): Handle new OpenMP target
	EXEC_OMP_* codes and new clauses.
	* gfortran.h (gfc_statement): Add ST_OMP_TARGET, ST_OMP_END_TARGET,
	ST_OMP_TARGET_DATA, ST_OMP_END_TARGET_DATA, ST_OMP_TARGET_UPDATE,
	ST_OMP_DECLARE_TARGET, ST_OMP_TEAMS, ST_OMP_END_TEAMS,
	ST_OMP_DISTRIBUTE, ST_OMP_END_DISTRIBUTE, ST_OMP_DISTRIBUTE_SIMD,
	ST_OMP_END_DISTRIBUTE_SIMD, ST_OMP_DISTRIBUTE_PARALLEL_DO,
	ST_OMP_END_DISTRIBUTE_PARALLEL_DO, ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
	ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD, ST_OMP_TARGET_TEAMS,
	ST_OMP_END_TARGET_TEAMS, ST_OMP_TEAMS_DISTRIBUTE,
	ST_OMP_END_TEAMS_DISTRIBUTE, ST_OMP_TEAMS_DISTRIBUTE_SIMD,
	ST_OMP_END_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TARGET_TEAMS_DISTRIBUTE,
	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE,
	ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD,
	ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
	ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO,
	ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
	ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
	ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
	ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD.
	(symbol_attribute): Add omp_declare_target field.
	(gfc_omp_depend_op, gfc_omp_map_op): New enums.
	(gfc_omp_namelist): Replace rop field with union
	containing reduction_op, depend_op and map_op.
	(OMP_LIST_DEPEND_IN, OMP_LIST_DEPEND_OUT): Remove.
	(OMP_LIST_DEPEND, OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM): New.
	(gfc_omp_clauses): Add num_teams, device, thread_limit,
	dist_sched_kind, dist_chunk_size fields.
	(gfc_common_head): Add omp_declare_target field.
	(gfc_exec_op): Add EXEC_OMP_TARGET, EXEC_OMP_TARGET_DATA,
	EXEC_OMP_TEAMS, EXEC_OMP_DISTRIBUTE, EXEC_OMP_DISTRIBUTE_SIMD,
	EXEC_OMP_DISTRIBUTE_PARALLEL_DO, EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
	EXEC_OMP_TARGET_TEAMS, EXEC_OMP_TEAMS_DISTRIBUTE,
	EXEC_OMP_TEAMS_DISTRIBUTE_SIMD, EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
	EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
	EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
	EXEC_OMP_TARGET_UPDATE.
	(gfc_add_omp_declare_target): New prototype.
	* match.h (gfc_match_omp_declare_target, gfc_match_omp_distribute,
	gfc_match_omp_distribute_parallel_do,
	gfc_match_omp_distribute_parallel_do_simd,
	gfc_match_omp_distribute_simd, gfc_match_omp_target,
	gfc_match_omp_target_data, gfc_match_omp_target_teams,
	gfc_match_omp_target_teams_distribute,
	gfc_match_omp_target_teams_distribute_parallel_do,
	gfc_match_omp_target_teams_distribute_parallel_do_simd,
	gfc_match_omp_target_teams_distribute_simd,
	gfc_match_omp_target_update, gfc_match_omp_teams,
	gfc_match_omp_teams_distribute,
	gfc_match_omp_teams_distribute_parallel_do,
	gfc_match_omp_teams_distribute_parallel_do_simd,
	gfc_match_omp_teams_distribute_simd): New prototypes.
	* module.c (ab_attribute): Add AB_OMP_DECLARE_TARGET.
	(attr_bits): Likewise.
	(mio_symbol_attribute): Handle omp_declare_target attribute.
	(gfc_free_omp_clauses): Free num_teams, device, thread_limit
	and dist_chunk_size expressions.
	(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): Use 1U instead of 1.
	(OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO, OMP_CLAUSE_FROM,
	OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT,
	OMP_CLAUSE_DIST_SCHEDULE): Define.
	(gfc_match_omp_clauses): Change mask parameter to unsigned int.
	Adjust for rop becoming u.reduction_op.  Disallow inbranch with
	notinbranch.  For depend clause, always create OMP_LIST_DEPEND
	and fill in u.depend_op.  Handle num_teams, device, map,
	to, from, thread_limit and dist_schedule clauses.
	(OMP_DECLARE_SIMD_CLAUSES): Or in OMP_CLAUSE_INBRANCH and
	OMP_CLAUSE_NOTINBRANCH.
	(OMP_TARGET_CLAUSES, OMP_TARGET_DATA_CLAUSES,
	OMP_TARGET_UPDATE_CLAUSES, OMP_TEAMS_CLAUSES,
	OMP_DISTRIBUTE_CLAUSES): Define.
	(match_omp): New function.
	(gfc_match_omp_do, gfc_match_omp_do_simd, gfc_match_omp_parallel,
	gfc_match_omp_parallel_do, gfc_match_omp_parallel_do_simd,
	gfc_match_omp_parallel_sections, gfc_match_omp_parallel_workshare,
	gfc_match_omp_sections, gfc_match_omp_simd, gfc_match_omp_single,
	gfc_match_omp_task): Rewritten using match_omp.
	(gfc_match_omp_threadprivate, gfc_match_omp_declare_reduction):
	Diagnose if the directives are followed by unexpected junk.
	(gfc_match_omp_distribute, gfc_match_omp_distribute_parallel_do,
	gfc_match_omp_distribute_parallel_do_simd,
	gfc_match_omp_distrbute_simd, gfc_match_omp_declare_target,
	gfc_match_omp_target, gfc_match_omp_target_data,
	gfc_match_omp_target_teams, gfc_match_omp_target_teams_distribute,
	gfc_match_omp_target_teams_distribute_parallel_do,
	gfc_match_omp_target_teams_distribute_parallel_do_simd,
	gfc_match_omp_target_teams_distrbute_simd, gfc_match_omp_target_update,
	gfc_match_omp_teams, gfc_match_omp_teams_distribute,
	gfc_match_omp_teams_distribute_parallel_do,
	gfc_match_omp_teams_distribute_parallel_do_simd,
	gfc_match_omp_teams_distrbute_simd): New functions.
	* openmp.c (resolve_omp_clauses): Adjust for
	OMP_LIST_DEPEND_{IN,OUT} being changed to OMP_LIST_DEPEND.  Handle
	OMP_LIST_MAP, OMP_LIST_FROM, OMP_LIST_TO, num_teams, device,
	dist_chunk_size and thread_limit.
	(gfc_resolve_omp_parallel_blocks): Only put sharing clauses into
	ctx.sharing_clauses.  Call gfc_resolve_omp_do_blocks for various
	new EXEC_OMP_* codes.
	(resolve_omp_do): Handle various new EXEC_OMP_* codes.
	(gfc_resolve_omp_directive): Likewise.
	(gfc_resolve_omp_declare_simd): Add missing space to diagnostics.
	* parse.c (decode_omp_directive): Handle parsing of OpenMP 4.0
	offloading related directives.
	(case_executable): Add ST_OMP_TARGET_UPDATE.
	(case_exec_markers): Add ST_OMP_TARGET*, ST_OMP_TEAMS*,
	ST_OMP_DISTRIBUTE*.
	(case_decl): Add ST_OMP_DECLARE_TARGET.
	(gfc_ascii_statement): Handle new ST_OMP_* codes.
	(parse_omp_do): Handle various new ST_OMP_* codes.
	(parse_executable): Likewise.
	* resolve.c (gfc_resolve_blocks): Handle various new EXEC_OMP_*
	codes.
	(resolve_code): Likewise.
	(resolve_symbol): Change that !$OMP DECLARE TARGET variables
	are saved.
	* st.c (gfc_free_statement): Handle various new EXEC_OMP_* codes.
	* symbol.c (check_conflict): Check omp_declare_target conflicts.
	(gfc_add_omp_declare_target): New function.
	(gfc_copy_attr): Copy omp_declare_target.
	* trans.c (trans_code): Handle various new EXEC_OMP_* codes.
	* trans-common.c (build_common_decl): Add "omp declare target"
	attribute if needed.
	* trans-decl.c (add_attributes_to_decl): Likewise.
	* trans.h (gfc_omp_finish_clause): New prototype.
	* trans-openmp.c (gfc_omp_finish_clause): New function.
	(gfc_trans_omp_reduction_list): Adjust for rop being renamed
	to u.reduction_op.
	(gfc_trans_omp_clauses): Adjust for OMP_LIST_DEPEND_{IN,OUT}
	change to OMP_LIST_DEPEND and fix up depend handling.
	Handle OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM, num_teams,
	thread_limit, device, dist_chunk_size and dist_sched_kind.
	(gfc_trans_omp_do): Handle EXEC_OMP_DISTRIBUTE.
	(GFC_OMP_SPLIT_DISTRIBUTE, GFC_OMP_SPLIT_TEAMS,
	GFC_OMP_SPLIT_TARGET, GFC_OMP_SPLIT_NUM, GFC_OMP_MASK_DISTRIBUTE,
	GFC_OMP_MASK_TEAMS, GFC_OMP_MASK_TARGET, GFC_OMP_MASK_NUM): New.
	(gfc_split_omp_clauses): Handle splitting of clauses for new
	EXEC_OMP_* codes.
	(gfc_trans_omp_do_simd): Add pblock argument, adjust for being
	callable for combined constructs.
	(gfc_trans_omp_parallel_do, gfc_trans_omp_parallel_do_simd): Likewise.
	(gfc_trans_omp_distribute, gfc_trans_omp_teams,
	gfc_trans_omp_target, gfc_trans_omp_target_data,
	gfc_trans_omp_target_update): New functions.
	(gfc_trans_omp_directive): Adjust gfc_trans_omp_* callers, handle
	new EXEC_OMP_* codes.
gcc/testsuite/
	* gfortran.dg/gomp/declare-simd-1.f90: New test.
	* gfortran.dg/gomp/depend-1.f90: New test.
	* gfortran.dg/gomp/target1.f90: New test.
	* gfortran.dg/gomp/target2.f90: New test.
	* gfortran.dg/gomp/target3.f90: New test.
	* gfortran.dg/gomp/udr4.f90: Adjust expected diagnostics.
	* gfortran.dg/openmp-define-3.f90: Expect _OPENMP 201307 instead of
	201107.
libgomp/
	* omp_lib.f90.in (openmp_version): Set to 201307.
	* omp_lib.h.in (openmp_version): Likewise.
	* testsuite/libgomp.c/target-8.c: New test.
	* testsuite/libgomp.fortran/declare-simd-1.f90: Add notinbranch
	and inbranch clauses.
	* testsuite/libgomp.fortran/depend-3.f90: New test.
	* testsuite/libgomp.fortran/openmp_version-1.f: Adjust for new
	openmp_version.
	* testsuite/libgomp.fortran/openmp_version-2.f90: Likewise.
	* testsuite/libgomp.fortran/target1.f90: New test.
	* testsuite/libgomp.fortran/target2.f90: New test.
	* testsuite/libgomp.fortran/target3.f90: New test.
	* testsuite/libgomp.fortran/target4.f90: New test.
	* testsuite/libgomp.fortran/target5.f90: New test.
	* testsuite/libgomp.fortran/target6.f90: New test.
	* testsuite/libgomp.fortran/target7.f90: New test.

From-SVN: r211768
2014-06-18 09:16:12 +02:00
Jakub Jelinek 92d28cbb59 re PR fortran/60928 (gfortran issue with allocatable components and OpenMP)
PR fortran/60928
	* omp-low.c (lower_rec_input_clauses) <case OMP_CLAUSE_LASTPRIVATE>:
	Set lastprivate_firstprivate even if omp_private_outer_ref
	langhook returns true.
	<case OMP_CLAUSE_REDUCTION>: When calling omp_clause_default_ctor
	langhook, call unshare_expr on new_var and call
	build_outer_var_ref to get the last argument.
gcc/c-family/
	* c-pragma.c (omp_pragmas_simd): Move PRAGMA_OMP_TASK...
	(omp_pragmas): ... back here.
gcc/fortran/
	* f95-lang.c (gfc_init_builtin_functions): Handle -fopenmp-simd
	like -fopenmp.
	* openmp.c (resolve_omp_clauses): Remove allocatable components
	diagnostics.  Add associate-name and intent(in) pointer
	diagnostics for various clauses, diagnose procedure pointers in
	reduction clause.
	* parse.c (match_word_omp_simd): New function.
	(matchs, matcho): New macros.
	(decode_omp_directive): Change match macros to either matchs
	or matcho.  Handle -fopenmp-simd.
	(next_free, next_fixed): Handle -fopenmp-simd like -fopenmp.
	* scanner.c (skip_free_comments, skip_fixed_comments, include_line):
	Likewise.
	* trans-array.c (get_full_array_size): Rename to...
	(gfc_full_array_size): ... this.  No longer static.
	(duplicate_allocatable): Adjust caller.  Add NO_MEMCPY argument
	and handle it.
	(gfc_duplicate_allocatable, gfc_copy_allocatable_data): Adjust
	duplicate_allocatable callers.
	(gfc_duplicate_allocatable_nocopy): New function.
	(structure_alloc_comps): Adjust g*_full_array_size and
	duplicate_allocatable caller.
	* trans-array.h (gfc_full_array_size,
	gfc_duplicate_allocatable_nocopy): New prototypes.
	* trans-common.c (create_common): Call gfc_finish_decl_attrs.
	* trans-decl.c (gfc_finish_decl_attrs): New function.
	(gfc_finish_var_decl, create_function_arglist,
	gfc_get_fake_result_decl): Call it.
	(gfc_allocate_lang_decl): If DECL_LANG_SPECIFIC is already allocated,
	don't allocate it again.
	(gfc_get_symbol_decl): Set GFC_DECL_ASSOCIATE_VAR_P on
	associate-names.
	* trans.h (gfc_finish_decl_attrs): New prototype.
	(struct lang_decl): Add scalar_allocatable and scalar_pointer
	bitfields.
	(GFC_DECL_SCALAR_ALLOCATABLE, GFC_DECL_SCALAR_POINTER,
	GFC_DECL_GET_SCALAR_ALLOCATABLE, GFC_DECL_GET_SCALAR_POINTER,
	GFC_DECL_ASSOCIATE_VAR_P): Define.
	(GFC_POINTER_TYPE_P): Remove.
	* trans-openmp.c (gfc_omp_privatize_by_reference): Don't check
	GFC_POINTER_TYPE_P, instead test GFC_DECL_GET_SCALAR_ALLOCATABLE,
	GFC_DECL_GET_SCALAR_POINTER or GFC_DECL_CRAY_POINTEE on decl.
	(gfc_omp_predetermined_sharing): Associate-names are predetermined.
	(enum walk_alloc_comps): New.
	(gfc_has_alloc_comps, gfc_omp_unshare_expr_r, gfc_omp_unshare_expr,
	gfc_walk_alloc_comps): New functions.
	(gfc_omp_private_outer_ref): Return true for scalar allocatables or
	decls with allocatable components.
	(gfc_omp_clause_default_ctor, gfc_omp_clause_copy_ctor,
	gfc_omp_clause_assign_op, gfc_omp_clause_dtor): Fix up handling of
	allocatables, handle also OMP_CLAUSE_REDUCTION, handle scalar
	allocatables and decls with allocatable components.
	(gfc_trans_omp_array_reduction_or_udr): Don't handle allocatable
	arrays here.
	(gfc_trans_omp_reduction_list): Call
	gfc_trans_omp_array_reduction_or_udr even for allocatable scalars.
	(gfc_trans_omp_do_simd): If -fno-openmp, just expand it as OMP_SIMD.
	(gfc_trans_omp_parallel_do_simd): Likewise.
	* trans-types.c (gfc_sym_type): Don't set GFC_POINTER_TYPE_P.
	(gfc_get_derived_type): Call gfc_finish_decl_attrs.
gcc/testsuite/
	* gfortran.dg/gomp/allocatable_components_1.f90: Remove dg-error
	directives.
	* gfortran.dg/gomp/associate1.f90: New test.
	* gfortran.dg/gomp/intentin1.f90: New test.
	* gfortran.dg/gomp/openmp-simd-1.f90: New test.
	* gfortran.dg/gomp/openmp-simd-2.f90: New test.
	* gfortran.dg/gomp/openmp-simd-3.f90: New test.
	* gfortran.dg/gomp/proc_ptr_2.f90: New test.
libgomp/
	* testsuite/libgomp.fortran/allocatable9.f90: New test.
	* testsuite/libgomp.fortran/allocatable10.f90: New test.
	* testsuite/libgomp.fortran/allocatable11.f90: New test.
	* testsuite/libgomp.fortran/allocatable12.f90: New test.
	* testsuite/libgomp.fortran/alloc-comp-1.f90: New test.
	* testsuite/libgomp.fortran/alloc-comp-2.f90: New test.
	* testsuite/libgomp.fortran/alloc-comp-3.f90: New test.
	* testsuite/libgomp.fortran/associate1.f90: New test.
	* testsuite/libgomp.fortran/associate2.f90: New test.
	* testsuite/libgomp.fortran/procptr1.f90: New test.

From-SVN: r211397
2014-06-10 08:05:22 +02:00
Jakub Jelinek 5f23671d3f dump-parse-tree.c (show_omp_namelist): Dump reduction id in each list item.
gcc/fortran/
	* dump-parse-tree.c (show_omp_namelist): Dump reduction
	id in each list item.
	(show_omp_node): Only handle OMP_LIST_REDUCTION, not
	OMP_LIST_REDUCTION_FIRST .. OMP_LIST_REDUCTION_LAST.  Don't
	dump reduction id here.
	* frontend-passes.c (dummy_code_callback): Renamed to...
	(gfc_dummy_code_callback): ... this.  No longer static.
	(optimize_reduction): Use gfc_dummy_code_callback instead of
	dummy_code_callback.
	* gfortran.h (gfc_statement): Add ST_OMP_DECLARE_REDUCTION.
	(symbol_attribute): Add omp_udr_artificial_var bitfield.
	(gfc_omp_reduction_op): New enum.
	(gfc_omp_namelist): Add rop and udr fields.
	(OMP_LIST_PLUS, OMP_LIST_REDUCTION_FIRST, OMP_LIST_MULT,
	OMP_LIST_SUB, OMP_LIST_AND, OMP_LIST_OR, OMP_LIST_EQV,
	OMP_LIST_NEQV, OMP_LIST_MAX, OMP_LIST_MIN, OMP_LIST_IAND,
	OMP_LIST_IOR, OMP_LIST_IEOR, OMP_LIST_REDUCTION_LAST): Removed.
	(OMP_LIST_REDUCTION): New.
	(gfc_omp_udr): New type.
	(gfc_get_omp_udr): Define.
	(gfc_symtree): Add n.omp_udr field.
	(gfc_namespace): Add omp_udr_root field, add omp_udr_ns bitfield.
	(gfc_free_omp_udr, gfc_omp_udr_find, gfc_resolve_omp_udrs,
	gfc_dummy_code_callback): New prototypes.
	* match.h (gfc_match_omp_declare_reduction): New prototype.
	* module.c (MOD_VERSION): Increase to 13.
	(omp_declare_reduction_stmt): New array.
	(mio_omp_udr_expr, write_omp_udr, write_omp_udrs, load_omp_udrs):
	New functions.
	(read_module): Read OpenMP user defined reductions.
	(write_module): Write OpenMP user defined reductions.
	* openmp.c: Include arith.h.
	(gfc_free_omp_udr, gfc_find_omp_udr): New functions.
	(gfc_match_omp_clauses): Handle user defined reductions.
	Store reduction kind into gfc_omp_namelist instead of using
	several OMP_LIST_* entries.
	(match_udr_expr, gfc_omp_udr_predef, gfc_omp_udr_find,
	gfc_match_omp_declare_reduction): New functions.
	(resolve_omp_clauses): Adjust for reduction clauses being only
	in OMP_LIST_REDUCTION list.  Diagnose missing UDRs.
	(struct omp_udr_callback_data): New type.
	(omp_udr_callback, gfc_resolve_omp_udr, gfc_resolve_omp_udrs): New
	functions.
	* parse.c (decode_omp_directive): Handle !$omp declare reduction.
	(case_decl): Add ST_OMP_DECLARE_REDUCTION.
	(gfc_ascii_statement): Print ST_OMP_DECLARE_REDUCTION.
	* resolve.c (resolve_fl_variable): Allow len=: or len=* on
	sym->attr.omp_udr_artificial_var symbols.
	(resolve_types): Call gfc_resolve_omp_udrs.
	* symbol.c (gfc_get_uop): If gfc_current_ns->omp_udr_ns,
	use parent ns instead of gfc_current_ns.
	(gfc_get_sym_tree): Don't insert symbols into
	namespaces with omp_udr_ns set.
	(free_omp_udr_tree): New function.
	(gfc_free_namespace): Call it.
	* trans-openmp.c (struct omp_udr_find_orig_data): New type.
	(omp_udr_find_orig, gfc_trans_omp_udr_expr): New functions.
	(gfc_trans_omp_array_reduction): Renamed to...
	(gfc_trans_omp_array_reduction_or_udr): ... this.  Remove SYM
	argument, instead pass gfc_omp_namelist pointer N.  Handle
	user defined reductions.
	(gfc_trans_omp_reduction_list): Remove REDUCTION_CODE argument.
	Handle user defined reductions and reduction ops in gfc_omp_namelist.
	(gfc_trans_omp_clauses): Adjust for just a single OMP_LIST_REDUCTION
	list.
	(gfc_split_omp_clauses): Likewise.
gcc/testsuite/
	* gfortran.dg/gomp/allocatable_components_1.f90: Adjust for
	reduction clause diagnostic changes.
	* gfortran.dg/gomp/appendix-a/a.31.3.f90: Likewise.
	* gfortran.dg/gomp/reduction1.f90: Likewise.
	* gfortran.dg/gomp/reduction3.f90: Likewise.
	* gfortran.dg/gomp/udr1.f90: New test.
	* gfortran.dg/gomp/udr2.f90: New test.
	* gfortran.dg/gomp/udr3.f90: New test.
	* gfortran.dg/gomp/udr4.f90: New test.
	* gfortran.dg/gomp/udr5.f90: New test.
	* gfortran.dg/gomp/udr6.f90: New test.
	* gfortran.dg/gomp/udr7.f90: New test.
libgomp/
	* testsuite/libgomp.fortran/simd1.f90: New test.
	* testsuite/libgomp.fortran/udr1.f90: New test.
	* testsuite/libgomp.fortran/udr2.f90: New test.
	* testsuite/libgomp.fortran/udr3.f90: New test.
	* testsuite/libgomp.fortran/udr4.f90: New test.
	* testsuite/libgomp.fortran/udr5.f90: New test.
	* testsuite/libgomp.fortran/udr6.f90: New test.
	* testsuite/libgomp.fortran/udr7.f90: New test.
	* testsuite/libgomp.fortran/udr8.f90: New test.
	* testsuite/libgomp.fortran/udr9.f90: New test.
	* testsuite/libgomp.fortran/udr10.f90: New test.
	* testsuite/libgomp.fortran/udr11.f90: New test.

From-SVN: r211303
2014-06-06 09:24:38 +02:00
Uros Bizjak 0389fbb55b declare-simd-1.f90: Require vect_simd_clones effective target.
* testsuite/libgomp.fortran/declare-simd-1.f90: Require
	vect_simd_clones effective target.
	* testsuite/libgomp.fortran/declare-simd-2.f90: Ditto.

From-SVN: r210961
2014-05-27 11:14:53 +02:00
Jakub Jelinek decaaec811 re PR middle-end/61252 (Invalid code produced for omp simd reduction(min:var) where var is reference)
PR middle-end/61252
	* omp-low.c (handle_simd_reference): New function.
	(lower_rec_input_clauses): Use it.  Defer adding reference
	initialization even for reduction without placeholder if in simd,
	handle it properly later on.

	* testsuite/libgomp.c++/simd-9.C: New test.

From-SVN: r210679
2014-05-21 10:04:03 +02:00
Jakub Jelinek dd2fc5256e tree.h (OMP_CLAUSE_LINEAR_STMT): Define.
* tree.h (OMP_CLAUSE_LINEAR_STMT): Define.
	* tree.c (omp_clause_num_ops): Increase OMP_CLAUSE_LINEAR
	number of operands to 3.
	(walk_tree_1): Walk all operands of OMP_CLAUSE_LINEAR.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Handle OMP_CLAUSE_DEPEND.
	* gimplify.c (gimplify_scan_omp_clauses): Handle
	OMP_CLAUSE_LINEAR_STMT.
	* omp-low.c (lower_rec_input_clauses): Fix typo.
	(maybe_add_implicit_barrier_cancel, lower_omp_1): Add
	cast between Fortran boolean_type_node and C _Bool if
	needed.
gcc/fortran/
	* gfortran.h (gfc_statement): Add ST_OMP_CANCEL,
	ST_OMP_CANCELLATION_POINT, ST_OMP_TASKGROUP, ST_OMP_END_TASKGROUP,
	ST_OMP_SIMD, ST_OMP_END_SIMD, ST_OMP_DO_SIMD, ST_OMP_END_DO_SIMD,
	ST_OMP_PARALLEL_DO_SIMD, ST_OMP_END_PARALLEL_DO_SIMD and
	ST_OMP_DECLARE_SIMD.
	(gfc_omp_namelist): New typedef.
	(gfc_get_omp_namelist): Define.
	(OMP_LIST_UNIFORM, OMP_LIST_ALIGNED, OMP_LIST_LINEAR,
	OMP_LIST_DEPEND_IN, OMP_LIST_DEPEND_OUT): New clause list kinds.
	(gfc_omp_proc_bind_kind, gfc_omp_cancel_kind): New enums.
	(gfc_omp_clauses): Change type of lists to gfc_omp_namelist *.
	Add inbranch, notinbranch, cancel, proc_bind, safelen_expr and
	simdlen_expr fields.
	(gfc_omp_declare_simd): New typedef.
	(gfc_get_omp_declare_simd): Define.
	(gfc_namespace): Add omp_declare_simd field.
	(gfc_exec_op): Add EXEC_OMP_CANCEL, EXEC_OMP_CANCELLATION_POINT,
	EXEC_OMP_TASKGROUP, EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD and
	EXEC_OMP_PARALLEL_DO_SIMD.
	(gfc_omp_atomic_op): Add GFC_OMP_ATOMIC_MASK, GFC_OMP_ATOMIC_SEQ_CST
	and GFC_OMP_ATOMIC_SWAP.
	(gfc_code): Change type of omp_namelist field to gfc_omp_namelist *.
	(gfc_free_omp_namelist, gfc_free_omp_declare_simd,
	gfc_free_omp_declare_simd_list, gfc_resolve_omp_declare_simd): New
	prototypes.
	* trans-stmt.h (gfc_trans_omp_declare_simd): New prototype.
	* symbol.c (gfc_free_namespace): Call gfc_free_omp_declare_simd.
	* openmp.c (gfc_free_omp_clauses): Free safelen_expr and
	simdlen_expr.  Use gfc_free_omp_namelist instead of
	gfc_free_namelist.
	(gfc_free_omp_declare_simd, gfc_free_omp_declare_simd_list): New
	functions.
	(gfc_match_omp_variable_list): Add end_colon, headp and
	allow_sections arguments.  Handle parsing of array sections.
	Use *omp_namelist* instead of *namelist* data structure and
	functions/macros.  Allow termination at : character.
	(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): Define.
	(gfc_match_omp_clauses): Change first and needs_space variables
	into arguments with default values.  Parse inbranch, notinbranch,
	proc_bind, safelen, simdlen, uniform, linear, aligned and
	depend clauses.
	(OMP_PARALLEL_CLAUSES): Add OMP_CLAUSE_PROC_BIND.
	(OMP_DECLARE_SIMD_CLAUSES, OMP_SIMD_CLAUSES): Define.
	(OMP_TASK_CLAUSES): Add OMP_CLAUSE_DEPEND.
	(gfc_match_omp_do_simd): New function.
	(gfc_match_omp_flush): Use *omp_namelist* instead of *namelist*
	data structure and functions/macros.
	(gfc_match_omp_simd, gfc_match_omp_declare_simd,
	gfc_match_omp_parallel_do_simd): New functions.
	(gfc_match_omp_atomic): Handle seq_cst clause.  Handle atomic swap.
	(gfc_match_omp_taskgroup, gfc_match_omp_cancel_kind,
	gfc_match_omp_cancel, gfc_match_omp_cancellation_point): New
	functions.
	(resolve_omp_clauses): Add where, omp_clauses and ns arguments.
	Use *omp_namelist* instead of *namelist* data structure and
	functions/macros.  Resolve uniform, aligned, linear, depend,
	safelen and simdlen clauses.
	(resolve_omp_atomic): Adjust for GFC_OMP_ATOMIC_{MASK,SEQ_CST,SWAP}
	addition, recognize atomic swap.
	(gfc_resolve_omp_parallel_blocks): Use gfc_omp_namelist instead
	of gfc_namelist.  Handle EXEC_OMP_PARALLEL_DO_SIMD the same as
	EXEC_OMP_PARALLEL_DO.
	(gfc_resolve_do_iterator): Use *omp_namelist* instead of *namelist*
	data structure and functions/macros.
	(resolve_omp_do): Likewise.  Handle EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD.
	(gfc_resolve_omp_directive): Handle EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD and EXEC_OMP_CANCEL.  Adjust
	resolve_omp_clauses caller.
	(gfc_resolve_omp_declare_simd): New function.
	* parse.c (decode_omp_directive): Parse cancellation point, cancel,
	declare simd, end do simd, end simd, end parallel do simd,
	end taskgroup, parallel do simd, simd and taskgroup directives.
	(case_executable): Add ST_OMP_CANCEL and ST_OMP_CANCELLATION_POINT.
	(case_exec_markers): Add ST_OMP_TASKGROUP, case ST_OMP_SIMD,
	ST_OMP_DO_SIMD and ST_OMP_PARALLEL_DO_SIMD.
	(case_decl): Add ST_OMP_DECLARE_SIMD.
	(gfc_ascii_statement): Handle ST_OMP_CANCEL,
	ST_OMP_CANCELLATION_POINT, ST_OMP_TASKGROUP, ST_OMP_END_TASKGROUP,
	ST_OMP_SIMD, ST_OMP_END_SIMD, ST_OMP_DO_SIMD, ST_OMP_END_DO_SIMD,
	ST_OMP_PARALLEL_DO_SIMD, ST_OMP_END_PARALLEL_DO_SIMD and
	ST_OMP_DECLARE_SIMD.
	(parse_omp_do): Handle ST_OMP_SIMD, ST_OMP_DO_SIMD and
	ST_OMP_PARALLEL_DO_SIMD.
	(parse_omp_atomic): Adjust for GFC_OMP_ATOMIC_* additions.
	(parse_omp_structured_block): Handle ST_OMP_TASKGROUP and
	ST_OMP_PARALLEL_DO_SIMD.
	(parse_executable): Handle ST_OMP_SIMD, ST_OMP_DO_SIMD,
	ST_OMP_PARALLEL_DO_SIMD and ST_OMP_TASKGROUP.
	* trans-decl.c (gfc_get_extern_function_decl,
	gfc_create_function_decl): Call gfc_trans_omp_declare_simd if
	needed.
	* frontend-passes.c (gfc_code_walker): Handle EXEC_OMP_SIMD,
	EXEC_OMP_DO_SIMD and EXEC_OMP_PARALLEL_DO_SIMD.  Walk
	safelen_expr and simdlen_expr.  Walk expressions in gfc_omp_namelist
	of depend, aligned and linear clauses.
	* match.c (match_exit_cycle): Handle EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD
	and EXEC_OMP_PARALLEL_DO_SIMD.
	(gfc_free_omp_namelist): New function.
	* dump-parse-tree.c (show_namelist): Removed.
	(show_omp_namelist): New function.
	(show_omp_node): Handle OpenMP 4.0 additions.
	(show_code_node): Handle EXEC_OMP_CANCEL, EXEC_OMP_CANCELLATION_POINT,
	EXEC_OMP_DO_SIMD, EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and
	EXEC_OMP_TASKGROUP.
	* match.h (gfc_match_omp_cancel, gfc_match_omp_cancellation_point,
	gfc_match_omp_declare_simd, gfc_match_omp_do_simd,
	gfc_match_omp_parallel_do_simd, gfc_match_omp_simd,
	gfc_match_omp_taskgroup): New prototypes.
	* trans-openmp.c (gfc_trans_omp_variable): Add declare_simd
	argument, handle it.  Allow current_function_decl to be NULL.
	(gfc_trans_omp_variable_list): Add declare_simd argument, pass
	it through to gfc_trans_omp_variable and disregard whether
	sym is referenced if declare_simd is true.  Work on gfc_omp_namelist
	instead of gfc_namelist.
	(gfc_trans_omp_reduction_list): Work on gfc_omp_namelist instead of
	gfc_namelist.  Adjust gfc_trans_omp_variable caller.
	(gfc_trans_omp_clauses): Add declare_simd argument, pass it through
	to gfc_trans_omp_variable{,_list} callers.  Work on gfc_omp_namelist
	instead of gfc_namelist.  Handle inbranch, notinbranch, safelen,
	simdlen, depend, uniform, linear, proc_bind and aligned clauses.
	Handle cancel kind.
	(gfc_trans_omp_atomic): Handle seq_cst clause, handle atomic swap,
	adjust for GFC_OMP_ATOMIC_* changes.
	(gfc_trans_omp_cancel, gfc_trans_omp_cancellation_point): New
	functions.
	(gfc_trans_omp_do): Add op argument, handle simd translation into
	generic.
	(GFC_OMP_SPLIT_SIMD, GFC_OMP_SPLIT_DO, GFC_OMP_SPLIT_PARALLEL,
	GFC_OMP_SPLIT_NUM, GFC_OMP_MASK_SIMD, GFC_OMP_MASK_DO,
	GFC_OMP_MASK_PARALLEL): New.
	(gfc_split_omp_clauses, gfc_trans_omp_do_simd): New functions.
	(gfc_trans_omp_parallel_do): Rework to use gfc_split_omp_clauses.
	(gfc_trans_omp_parallel_do_simd, gfc_trans_omp_taskgroup): New
	functions.
	(gfc_trans_omp_directive): Handle EXEC_OMP_CANCEL,
	EXEC_OMP_CANCELLATION_POINT, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and EXEC_OMP_TASKGROUP.
	Adjust gfc_trans_omp_do caller.
	(gfc_trans_omp_declare_simd): New function.
	* st.c (gfc_free_statement): Handle EXEC_OMP_CANCEL,
	EXEC_OMP_CANCELLATION_POINT, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and EXEC_OMP_TASKGROUP.
	For EXEC_OMP_FLUSH call gfc_free_omp_namelist instead of
	gfc_free_namelist.
	* module.c (omp_declare_simd_clauses): New variable.
	(mio_omp_declare_simd): New function.
	(mio_symbol): Call it.
	* trans.c (trans_code): Handle EXEC_OMP_CANCEL,
	EXEC_OMP_CANCELLATION_POINT, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and EXEC_OMP_TASKGROUP.
	* resolve.c (gfc_resolve_blocks): Handle EXEC_OMP_DO_SIMD,  
	EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and EXEC_OMP_TASKGROUP.
	(resolve_code): Handle EXEC_OMP_CANCEL,
	EXEC_OMP_CANCELLATION_POINT, EXEC_OMP_DO_SIMD,
	EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_SIMD and EXEC_OMP_TASKGROUP.
	(resolve_types): Call gfc_resolve_omp_declare_simd.
gcc/testsuite/
	* gfortran.dg/gomp/affinity-1.f90: New test.
libgomp/
	* testsuite/libgomp.fortran/cancel-do-1.f90: New test.
	* testsuite/libgomp.fortran/cancel-do-2.f90: New test.
	* testsuite/libgomp.fortran/cancel-parallel-1.f90: New test.
	* testsuite/libgomp.fortran/cancel-parallel-3.f90: New test.
	* testsuite/libgomp.fortran/cancel-sections-1.f90: New test.
	* testsuite/libgomp.fortran/cancel-taskgroup-2.f90: New test.
	* testsuite/libgomp.fortran/declare-simd-1.f90: New test.
	* testsuite/libgomp.fortran/declare-simd-2.f90: New test.
	* testsuite/libgomp.fortran/declare-simd-3.f90: New test.
	* testsuite/libgomp.fortran/depend-1.f90: New test.
	* testsuite/libgomp.fortran/depend-2.f90: New test.
	* testsuite/libgomp.fortran/omp_atomic5.f90: New test.
	* testsuite/libgomp.fortran/simd1.f90: New test.
	* testsuite/libgomp.fortran/simd2.f90: New test.
	* testsuite/libgomp.fortran/simd3.f90: New test.
	* testsuite/libgomp.fortran/simd4.f90: New test.
	* testsuite/libgomp.fortran/taskgroup1.f90: New test.

From-SVN: r210313
2014-05-11 22:26:36 +02:00
Jakub Jelinek 95782571f3 gimplify.c (gimplify_adjust_omp_clauses_1): Handle GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE.
* gimplify.c (gimplify_adjust_omp_clauses_1): Handle
	GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE.
	(gimplify_adjust_omp_clauses): Simd region is never
	directly nested in combined parallel.  Instead, for linear
	with copyin/copyout, if in combined for simd loop, make decl
	firstprivate/lastprivate on OMP_FOR.
	* omp-low.c (expand_omp_for_generic, expand_omp_for_static_nochunk,
	expand_omp_for_static_chunk): When setting endvar, also set
	fd->loop.v to the same value.
libgomp/
	* testsuite/libgomp.c/simd-10.c: New test.
	* testsuite/libgomp.c/simd-11.c: New test.
	* testsuite/libgomp.c/simd-12.c: New test.
	* testsuite/libgomp.c/simd-13.c: New test.

From-SVN: r210009
2014-05-02 19:43:40 +02:00
Jakub Jelinek 42056eaced c-parser.c (c_parser_omp_atomic): Allow seq_cst before atomic-clause...
gcc/c/
	* c-parser.c (c_parser_omp_atomic): Allow seq_cst before
	atomic-clause, allow comma in between atomic-clause and
	seq_cst.
gcc/cp/
	* parser.c (cp_parser_omp_atomic): Allow seq_cst before
	atomic-clause, allow comma in between atomic-clause and
	seq_cst.
gcc/testsuite/
	* c-c++-common/gomp/atomic-16.c: Remove all dg-error directives.
	Replace load with read and store with write.
libgomp/
	* testsuite/libgomp.c++/atomic-14.C: Allow seq_cst and
	atomic type clauses in any order and optional comma in between.
	* testsuite/libgomp.c++/atomic-15.C: Likewise.
	* testsuite/libgomp.c/atomic-17.c: Likewise.

From-SVN: r209762
2014-04-24 23:20:28 +02:00