* 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.


git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@211768 138bc75d-0d04-0410-961f-82ee72b054a4

12 files changed, 409 insertions, 3 deletions

diff --git a/libgomp/testsuite/libgomp.c/target-8.c b/libgomp/testsuite/libgomp.c/target-8.c
new file mode 100644
index 00000000000..35084575324
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-8.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-options "-fopenmp" } */
+
+void
+foo (int *p)
+{
+  int i;
+  #pragma omp parallel
+  #pragma omp single
+  #pragma omp target teams distribute parallel for map(p[0:24])
+  for (i = 0; i < 24; i++)
+    p[i] = p[i] + 1;
+}
+
+int
+main ()
+{
+  int p[24], i;
+  for (i = 0; i < 24; i++)
+    p[i] = i;
+  foo (p);
+  for (i = 0; i < 24; i++)
+    if (p[i] != i + 1)
+      __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
index b141a3b3eda..5cd592c09db 100644
--- a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
@@ -6,7 +6,8 @@
 module declare_simd_1_mod
   contains
     real function foo (a, b, c)
-      !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5)
+      !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5) &
+      !$omp & notinbranch
       double precision, value :: a
       real, value :: c
       !$omp declare simd (foo)
@@ -22,6 +23,7 @@ end module declare_simd_1_mod
       real, value :: c
       real :: bar
       !$omp declare simd (bar) simdlen (4) linear (b : 2)
+      !$omp declare simd (bar) simdlen (16) inbranch
       double precision, value :: a
     end function bar
   end interface
diff --git a/libgomp/testsuite/libgomp.fortran/depend-3.f90 b/libgomp/testsuite/libgomp.fortran/depend-3.f90
new file mode 100644
index 00000000000..11be6410692
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/depend-3.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+
+  integer :: x(2, 3)
+  integer, allocatable :: z(:, :)
+  allocate (z(-2:3, 2:4))
+  call foo (x, z)
+contains
+  subroutine foo (x, z)
+    integer :: x(:, :), y
+    integer, allocatable :: z(:, :)
+    y = 1
+    !$omp parallel shared (x, y, z)
+      !$omp single
+        !$omp taskgroup
+          !$omp task depend(in: x)
+  	  if (y.ne.1) call abort
+          !$omp end task
+          !$omp task depend(out: x(1:2, 1:3))
+  	  y = 2
+          !$omp end task
+        !$omp end taskgroup
+        !$omp taskgroup
+          !$omp task depend(in: z)
+  	  if (y.ne.2) call abort
+          !$omp end task
+          !$omp task depend(out: z(-2:3, 2:4))
+  	  y = 3
+          !$omp end task
+        !$omp end taskgroup
+        !$omp taskgroup
+          !$omp task depend(in: x)
+  	  if (y.ne.3) call abort
+          !$omp end task
+          !$omp task depend(out: x(1:, 1:))
+  	  y = 4
+          !$omp end task
+        !$omp end taskgroup
+      !$omp end single
+    !$omp end parallel
+    if (y.ne.4) call abort
+  end subroutine
+end
diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
index aaa888189b1..be24adcca0c 100644
--- a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
+++ b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
@@ -4,6 +4,6 @@
       implicit none
       include "omp_lib.h"
 
-      if (openmp_version .ne. 201107) call abort;
+      if (openmp_version .ne. 201307) call abort;
 
       end program main
diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
index b2d1d261f27..62712c7d206 100644
--- a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
@@ -4,6 +4,6 @@ program main
   use omp_lib
   implicit none
 
-  if (openmp_version .ne. 201107) call abort;
+  if (openmp_version .ne. 201307) call abort;
 
 end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target1.f90 b/libgomp/testsuite/libgomp.fortran/target1.f90
new file mode 100644
index 00000000000..c70daace497
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target1.f90
@@ -0,0 +1,58 @@
+! { dg-do run }
+
+module target1
+contains
+  subroutine foo (p, v, w, n)
+    double precision, pointer :: p(:), v(:), w(:)
+    double precision :: q(n)
+    integer :: i, n
+    !$omp target if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+        q(i) = p(i)
+      end do
+    !$omp end target
+    if (any (p /= q)) call abort
+    do i = 1, n
+      if (p(i) /= i * iand (i, 63)) call abort
+    end do
+    !$omp target data if (n > 256) map (to: v(1:n), w) map (from: p, q)
+    !$omp target if (n > 256)
+      do i = 1, n
+        p(i) = 1.0
+        q(i) = 2.0
+      end do
+    !$omp end target
+    !$omp target if (n > 256)
+      do i = 1, n
+        p(i) = p(i) + v(i) * w(i)
+        q(i) = q(i) + v(i) * w(i)
+      end do
+    !$omp end target
+    !$omp target if (n > 256)
+      !$omp teams distribute parallel do simd linear(i:1)
+      do i = 1, n
+        p(i) = p(i) + 2.0
+        q(i) = q(i) + 3.0
+      end do
+    !$omp end target
+    !$omp end target data
+    if (any (p + 2.0 /= q)) call abort
+  end subroutine
+end module target1
+  use target1, only : foo
+  integer :: n, i
+  double precision, pointer :: p(:), v(:), w(:)
+  n = 10000
+  allocate (p(n), v(n), w(n))
+  do i = 1, n
+    v(i) = i
+    w(i) = iand (i, 63)
+  end do
+  call foo (p, v, w, n)
+  do i = 1, n
+    if (p(i) /= i * iand (i, 63) + 3) call abort
+  end do
+  deallocate (p, v, w)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target2.f90 b/libgomp/testsuite/libgomp.fortran/target2.f90
new file mode 100644
index 00000000000..42f704f2bb3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target2.f90
@@ -0,0 +1,96 @@
+! { dg-do run }
+! { dg-options "-fopenmp -ffree-line-length-160" }
+
+module target2
+contains
+  subroutine foo (a, b, c, d, e, f, g, n, q)
+    integer :: n, q
+    integer :: a, b(3:n), c(5:), d(2:*), e(:,:)
+    integer, pointer :: f, g(:)
+    integer :: h, i(3:n)
+    integer, pointer :: j, k(:)
+    logical :: r
+    allocate (j, k(4:n))
+    h = 14
+    i = 15
+    j = 16
+    k = 17
+    !$omp target map (to: a, b, c, d(2:n+1), e, f, g, h, i, j, k, n) map (from: r)
+      r = a /= 7
+      r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (f /= 12)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (h /= 14)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (j /= 16)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: b(3:n), c(5:n+4), d(2:n+1), e(1:,:2), g(3:n), i(3:n), k(4:n), n) map (from: r)
+      r = (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: b(5:n-2), c(7:n), d(4:n-2), e(1:,2:), g(5:n-3), i(6:n-4), k(5:n-5), n) map (from: r)
+      r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    !$omp target map (to: b(q+5:n-2+q), c(q+7:q+n), d(q+4:q+n-2), e(1:q+2,2:q+2), g(5+q:n-3+q), &
+    !$omp & i(6+q:n-4+q), k(5+q:n-5+q), n) map (from: r)
+      r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: d(2:n+1), n)
+      r = a /= 7
+      r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (f /= 12)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (h /= 14)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (j /= 16)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+  end subroutine foo
+end module target2
+  use target2, only : foo
+  integer, parameter :: n = 15, q = 0
+  integer :: a, b(2:n-1), c(n), d(n), e(3:4, 3:4)
+  integer, pointer :: f, g(:)
+  allocate (f, g(3:n))
+  a = 7
+  b = 8
+  c = 9
+  d = 10
+  e = 11
+  f = 12
+  g = 13
+  call foo (a, b, c, d, e, f, g, n, q)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target3.f90 b/libgomp/testsuite/libgomp.fortran/target3.f90
new file mode 100644
index 00000000000..1f197acdef7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target3.f90
@@ -0,0 +1,29 @@
+! { dg-do run }
+
+module target3
+contains
+  subroutine foo (f, g)
+    integer :: n
+    integer, pointer :: f, g(:)
+    integer, pointer :: j, k(:)
+    logical :: r
+    nullify (j)
+    k => null ()
+    !$omp target map (tofrom: f, g, j, k) map (from: r)
+      r = associated (f) .or. associated (g)
+      r = r .or. associated (j) .or. associated (k)
+    !$omp end target
+    if (r) call abort
+    !$omp target
+      r = associated (f) .or. associated (g)
+      r = r .or. associated (j) .or. associated (k)
+    !$omp end target
+    if (r) call abort
+  end subroutine foo
+end module target3
+  use target3, only : foo
+  integer, pointer :: f, g(:)
+  f => null ()
+  nullify (g)
+  call foo (f, g)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target4.f90 b/libgomp/testsuite/libgomp.fortran/target4.f90
new file mode 100644
index 00000000000..aa2f0a5ac19
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target4.f90
@@ -0,0 +1,48 @@
+! { dg-do run }
+
+module target4
+contains
+  subroutine foo (a,m,n)
+    integer :: m,n,i,j
+    double precision :: a(m, n), t
+    !$omp target data map(a) map(to: m, n)
+    do i=1,n
+      t = 0.0d0
+      !$omp target
+        !$omp parallel do reduction(+:t)
+          do j=1,m
+            t = t + a(j,i) * a(j,i)
+          end do
+      !$omp end target
+      t = 2.0d0 * t
+      !$omp target
+        !$omp parallel do
+          do j=1,m
+            a(j,i) = a(j,i) * t
+          end do
+      !$omp end target
+    end do
+    !$omp end target data
+  end subroutine foo
+end module target4
+  use target4, only : foo
+  integer :: i, j
+  double precision :: a(8, 9), res(8, 9)
+  do i = 1, 8
+    do j = 1, 9
+      a(i, j) = i + j
+    end do
+  end do
+  call foo (a, 8, 9)
+  res = reshape ((/ 1136.0d0, 1704.0d0, 2272.0d0, 2840.0d0, 3408.0d0, 3976.0d0, &
+&   4544.0d0, 5112.0d0, 2280.0d0, 3040.0d0, 3800.0d0, 4560.0d0, 5320.0d0, 6080.0d0, &
+&   6840.0d0, 7600.0d0, 3936.0d0, 4920.0d0, 5904.0d0, 6888.0d0, 7872.0d0, 8856.0d0, &
+&   9840.0d0, 10824.0d0, 6200.0d0, 7440.0d0, 8680.0d0, 9920.0d0, 11160.0d0, 12400.0d0, &
+&   13640.0d0, 14880.0d0, 9168.0d0, 10696.0d0, 12224.0d0, 13752.0d0, 15280.0d0, 16808.0d0, &
+&   18336.0d0, 19864.0d0, 12936.0d0, 14784.0d0, 16632.0d0, 18480.0d0, 20328.0d0, 22176.0d0, &
+&   24024.0d0, 25872.0d0, 17600.0d0, 19800.0d0, 22000.0d0, 24200.0d0, 26400.0d0, 28600.0d0, &
+&   30800.0d0, 33000.0d0, 23256.0d0, 25840.0d0, 28424.0d0, 31008.0d0, 33592.0d0, 36176.0d0, &
+&   38760.0d0, 41344.0d0, 30000.0d0, 33000.0d0, 36000.0d0, 39000.0d0, 42000.0d0, 45000.0d0, &
+&   48000.0d0, 51000.0d0 /), (/ 8, 9 /))
+  if (any (a /= res)) call abort
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target5.f90 b/libgomp/testsuite/libgomp.fortran/target5.f90
new file mode 100644
index 00000000000..c46faf226f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target5.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+  integer :: r
+  r = 0
+  call foo (r)
+  if (r /= 11) call abort
+contains
+  subroutine foo (r)
+    integer :: i, r
+    !$omp parallel
+    !$omp single
+    !$omp target teams distribute parallel do reduction (+: r)
+      do i = 1, 10
+        r = r + 1
+      end do
+      r = r + 1
+    !$omp end single
+    !$omp end parallel
+  end subroutine
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target6.f90 b/libgomp/testsuite/libgomp.fortran/target6.f90
new file mode 100644
index 00000000000..13f5a52edd2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target6.f90
@@ -0,0 +1,50 @@
+! { dg-do run }
+
+module target6
+contains
+  subroutine foo (p, v, w, n)
+    double precision, pointer :: p(:), v(:), w(:)
+    double precision :: q(n)
+    integer :: i, n
+    !$omp target data if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+    !$omp target if (n > 256)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+        q(i) = p(i)
+      end do
+    !$omp end target
+    !$omp target update if (n > 256) from (p)
+    do i = 1, n
+      if (p(i) /= i * iand (i, 63)) call abort
+      v(i) = v(i) + 1
+    end do
+    !$omp target update if (n > 256) to (v(1:n))
+    !$omp target if (n > 256)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+      end do
+    !$omp end target
+    !$omp end target data
+    do i = 1, n
+      if (q(i) /= (v(i) - 1) * w(i)) call abort
+      if (p(i) /= q(i) + w(i)) call abort
+    end do
+  end subroutine
+end module target6
+  use target6, only : foo
+  integer :: n, i
+  double precision, pointer :: p(:), v(:), w(:)
+  n = 10000
+  allocate (p(n), v(n), w(n))
+  do i = 1, n
+    v(i) = i
+    w(i) = iand (i, 63)
+  end do
+  call foo (p, v, w, n)
+  do i = 1, n
+    if (p(i) /= (i + 1) * iand (i, 63)) call abort
+  end do
+  deallocate (p, v, w)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target7.f90 b/libgomp/testsuite/libgomp.fortran/target7.f90
new file mode 100644
index 00000000000..4af0ee371bd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target7.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+  interface
+    real function foo (x)
+      !$omp declare target
+      real, intent(in) :: x
+    end function foo
+  end interface
+  integer, parameter :: n = 1000
+  integer, parameter :: c = 100
+  integer :: i, j
+  real :: a(n)
+  do i = 1, n
+    a(i) = i
+  end do
+  do i = 1, n, c
+    !$omp task shared(a)
+      !$omp target map(a(i:i+c-1))
+        !$omp parallel do
+          do j = i, i + c - 1
+            a(j) = foo (a(j))
+          end do
+      !$omp end target
+    !$omp end task
+  end do
+  do i = 1, n
+    if (a(i) /= i + 1) call abort
+  end do
+end
+real function foo (x)
+  !$omp declare target
+  real, intent(in) :: x
+  foo = x + 1
+end function foo