--- gcc/gimplify.c.jj	2015-10-22 13:41:23.214882261 +0200
+++ gcc/gimplify.c	2015-10-22 18:01:45.125588093 +0200
@@ -5849,9 +5849,10 @@ omp_default_clause (struct gimplify_omp_
 	    {
 	      splay_tree_node n2;
 
-	      if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
-		continue;
 	      n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
+	      if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0
+		  && (n2 == NULL || (n2->value & GOVD_DATA_SHARE_CLASS) == 0))
+		continue;
 	      if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
 		{
 		  flags |= GOVD_FIRSTPRIVATE;
@@ -6143,10 +6144,12 @@ omp_check_private (struct gimplify_omp_c
 	  return true;
 	}
 
-      if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+      n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+
+      if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+	  && (n == NULL || (n->value & GOVD_DATA_SHARE_CLASS) == 0))
 	continue;
 
-      n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
       if (n != NULL)
 	{
 	  if ((n->value & GOVD_LOCAL) != 0
@@ -6177,12 +6180,12 @@ omp_no_lastprivate (struct gimplify_omp_
 	  if (!ctx->combined_loop)
 	    return false;
 	  if (ctx->distribute)
-	    return true;
+	    return lang_GNU_Fortran ();
 	  break;
 	case ORT_COMBINED_PARALLEL:
 	  break;
 	case ORT_COMBINED_TEAMS:
-	  return true;
+	  return lang_GNU_Fortran ();
 	default:
 	  return false;
 	}
@@ -6279,16 +6282,25 @@ gimplify_scan_omp_clauses (tree *list_p,
 	  else if (error_operand_p (decl))
 	    goto do_add;
 	  else if (outer_ctx
-		   && outer_ctx->region_type == ORT_COMBINED_PARALLEL
+		   && (outer_ctx->region_type == ORT_COMBINED_PARALLEL
+		       || outer_ctx->region_type == ORT_COMBINED_TEAMS)
 		   && splay_tree_lookup (outer_ctx->variables,
 					 (splay_tree_key) decl) == NULL)
-	    omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+	    {
+	      omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+	      if (outer_ctx->outer_context)
+		omp_notice_variable (outer_ctx->outer_context, decl, true);
+	    }
 	  else if (outer_ctx
 		   && (outer_ctx->region_type & ORT_TASK) != 0
 		   && outer_ctx->combined_loop
 		   && splay_tree_lookup (outer_ctx->variables,
 					 (splay_tree_key) decl) == NULL)
-	    omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+	    {
+	      omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+	      if (outer_ctx->outer_context)
+		omp_notice_variable (outer_ctx->outer_context, decl, true);
+	    }
 	  else if (outer_ctx
 		   && outer_ctx->region_type == ORT_WORKSHARE
 		   && outer_ctx->combined_loop
@@ -6302,8 +6314,14 @@ gimplify_scan_omp_clauses (tree *list_p,
 		      == ORT_COMBINED_PARALLEL)
 		  && splay_tree_lookup (outer_ctx->outer_context->variables,
 					(splay_tree_key) decl) == NULL)
-		omp_add_variable (outer_ctx->outer_context, decl,
-				  GOVD_SHARED | GOVD_SEEN);
+		{
+		  struct gimplify_omp_ctx *octx = outer_ctx->outer_context;
+		  omp_add_variable (octx, decl, GOVD_SHARED | GOVD_SEEN);
+		  if (octx->outer_context)
+		    omp_notice_variable (octx->outer_context, decl, true);
+		}
+	      else if (outer_ctx->outer_context)
+		omp_notice_variable (outer_ctx->outer_context, decl, true);
 	    }
 	  goto do_add;
 	case OMP_CLAUSE_REDUCTION:
@@ -6416,9 +6434,7 @@ gimplify_scan_omp_clauses (tree *list_p,
 		    {
 		      if (octx->outer_context
 			  && (octx->outer_context->region_type
-			      == ORT_COMBINED_PARALLEL
-			      || (octx->outer_context->region_type
-				  == ORT_COMBINED_TEAMS)))
+			      == ORT_COMBINED_PARALLEL))
 			octx = octx->outer_context;
 		      else if (omp_check_private (octx, decl, false))
 			break;
@@ -6433,8 +6449,15 @@ gimplify_scan_omp_clauses (tree *list_p,
 			   && octx == outer_ctx)
 		    flags = GOVD_SEEN | GOVD_SHARED;
 		  else if (octx
+			   && octx->region_type == ORT_COMBINED_TEAMS)
+		    flags = GOVD_SEEN | GOVD_SHARED;
+		  else if (octx
 			   && octx->region_type == ORT_COMBINED_TARGET)
-		    flags &= ~GOVD_LASTPRIVATE;
+		    {
+		      flags &= ~GOVD_LASTPRIVATE;
+		      if (flags == GOVD_SEEN)
+			break;
+		    }
 		  else
 		    break;
 		  splay_tree_node on
@@ -7289,6 +7312,15 @@ gimplify_adjust_omp_clauses (gimple_seq
 	      else
 		OMP_CLAUSE_CODE (c) = OMP_CLAUSE_PRIVATE;
 	    }
+	  else if (code == OMP_DISTRIBUTE
+		   && OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
+	    {
+	      remove = true;
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"same variable used in %<firstprivate%> and "
+			"%<lastprivate%> clauses on %<distribute%> "
+			"construct");
+	    }
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -7902,6 +7934,26 @@ gimplify_omp_for (tree *expr_p, gimple_s
 			  OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
 			  flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
 			}
+		      else
+			{
+			  struct gimplify_omp_ctx *octx = outer->outer_context;
+			  if (octx
+			      && octx->region_type == ORT_COMBINED_PARALLEL
+			      && octx->outer_context
+			      && (octx->outer_context->region_type
+				  == ORT_WORKSHARE)
+			      && octx->outer_context->combined_loop)
+			    {
+			      octx = octx->outer_context;
+			      n = splay_tree_lookup (octx->variables,
+						     (splay_tree_key)decl);
+			      if (n != NULL && (n->value & GOVD_LOCAL) != 0)
+				{
+				  OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
+				  flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
+				}
+			    }
+			}
 		    }
 		}
 
@@ -7936,7 +7988,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
 			{
 			  omp_add_variable (outer, decl,
 					    GOVD_LASTPRIVATE | GOVD_SEEN);
-			  if (outer->outer_context)
+			  if (outer->region_type == ORT_COMBINED_PARALLEL
+			      && outer->outer_context
+			      && (outer->outer_context->region_type
+				  == ORT_WORKSHARE)
+			      && outer->outer_context->combined_loop)
+			    {
+			      outer = outer->outer_context;
+			      n = splay_tree_lookup (outer->variables,
+						     (splay_tree_key)decl);
+			      if (omp_check_private (outer, decl, false))
+				outer = NULL;
+			      else if (n == NULL
+				       || ((n->value & GOVD_DATA_SHARE_CLASS)
+					   == 0))
+				omp_add_variable (outer, decl,
+						  GOVD_LASTPRIVATE
+						  | GOVD_SEEN);
+			      else
+				outer = NULL;
+			    }
+			  if (outer && outer->outer_context
+			      && (outer->outer_context->region_type
+				  == ORT_COMBINED_TEAMS))
+			    {
+			      outer = outer->outer_context;
+			      n = splay_tree_lookup (outer->variables,
+						     (splay_tree_key)decl);
+			      if (n == NULL
+				  || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+				omp_add_variable (outer, decl,
+						  GOVD_SHARED | GOVD_SEEN);
+			      else
+				outer = NULL;
+			    }
+			  if (outer && outer->outer_context)
 			    omp_notice_variable (outer->outer_context, decl,
 						 true);
 			}
@@ -7985,7 +8071,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
 			{
 			  omp_add_variable (outer, decl,
 					    GOVD_LASTPRIVATE | GOVD_SEEN);
-			  if (outer->outer_context)
+			  if (outer->region_type == ORT_COMBINED_PARALLEL
+			      && outer->outer_context
+			      && (outer->outer_context->region_type
+				  == ORT_WORKSHARE)
+			      && outer->outer_context->combined_loop)
+			    {
+			      outer = outer->outer_context;
+			      n = splay_tree_lookup (outer->variables,
+						     (splay_tree_key)decl);
+			      if (omp_check_private (outer, decl, false))
+				outer = NULL;
+			      else if (n == NULL
+				       || ((n->value & GOVD_DATA_SHARE_CLASS)
+					   == 0))
+				omp_add_variable (outer, decl,
+						  GOVD_LASTPRIVATE
+						  | GOVD_SEEN);
+			      else
+				outer = NULL;
+			    }
+			  if (outer && outer->outer_context
+			      && (outer->outer_context->region_type
+				  == ORT_COMBINED_TEAMS))
+			    {
+			      outer = outer->outer_context;
+			      n = splay_tree_lookup (outer->variables,
+						     (splay_tree_key)decl);
+			      if (n == NULL
+				  || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+				omp_add_variable (outer, decl,
+						  GOVD_SHARED | GOVD_SEEN);
+			      else
+				outer = NULL;
+			    }
+			  if (outer && outer->outer_context)
 			    omp_notice_variable (outer->outer_context, decl,
 						 true);
 			}
--- gcc/omp-low.c.jj	2015-10-22 14:11:23.488003543 +0200
+++ gcc/omp-low.c	2015-10-23 13:10:16.995777473 +0200
@@ -2646,12 +2646,15 @@ add_taskreg_looptemp_clauses (enum gf_ma
 	  && TREE_CODE (fd.loop.n2) != INTEGER_CST)
 	{
 	  count += fd.collapse - 1;
-	  /* For taskloop, if there are lastprivate clauses on the inner
+	  /* If there are lastprivate clauses on the inner
 	     GIMPLE_OMP_FOR, add one more temporaries for the total number
 	     of iterations (product of count1 ... countN-1).  */
-	  if (msk == GF_OMP_FOR_KIND_TASKLOOP
-	      && find_omp_clause (gimple_omp_for_clauses (for_stmt),
-				  OMP_CLAUSE_LASTPRIVATE))
+	  if (find_omp_clause (gimple_omp_for_clauses (for_stmt),
+			       OMP_CLAUSE_LASTPRIVATE))
+	    count++;
+	  else if (msk == GF_OMP_FOR_KIND_FOR
+		   && find_omp_clause (gimple_omp_parallel_clauses (stmt),
+				       OMP_CLAUSE_LASTPRIVATE))
 	    count++;
 	}
       for (i = 0; i < count; i++)
@@ -8648,6 +8651,30 @@ expand_omp_for_static_nochunk (struct om
 				OMP_CLAUSE__LOOPTEMP_);
       gcc_assert (innerc);
       endvar = OMP_CLAUSE_DECL (innerc);
+      if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+	  && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+	{
+	  int i;
+	  for (i = 1; i < fd->collapse; i++)
+	    {
+	      innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+					OMP_CLAUSE__LOOPTEMP_);
+	      gcc_assert (innerc);
+	    }
+	  innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+				    OMP_CLAUSE__LOOPTEMP_);
+	  if (innerc)
+	    {
+	      /* If needed (distribute parallel for with lastprivate),
+		 propagate down the total number of iterations.  */
+	      tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+				     fd->loop.n2);
+	      t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+					    GSI_CONTINUE_LINKING);
+	      assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+	      gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+	    }
+	}
     }
   t = fold_convert (itype, s0);
   t = fold_build2 (MULT_EXPR, itype, t, step);
@@ -9133,6 +9160,30 @@ expand_omp_for_static_chunk (struct omp_
 				OMP_CLAUSE__LOOPTEMP_);
       gcc_assert (innerc);
       endvar = OMP_CLAUSE_DECL (innerc);
+      if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+	  && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+	{
+	  int i;
+	  for (i = 1; i < fd->collapse; i++)
+	    {
+	      innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+					OMP_CLAUSE__LOOPTEMP_);
+	      gcc_assert (innerc);
+	    }
+	  innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+				    OMP_CLAUSE__LOOPTEMP_);
+	  if (innerc)
+	    {
+	      /* If needed (distribute parallel for with lastprivate),
+		 propagate down the total number of iterations.  */
+	      tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+				     fd->loop.n2);
+	      t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+					    GSI_CONTINUE_LINKING);
+	      assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+	      gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+	    }
+	}
     }
 
   t = fold_convert (itype, s0);
@@ -13456,26 +13507,36 @@ lower_omp_for_lastprivate (struct omp_fo
       && TREE_CODE (n2) != INTEGER_CST
       && gimple_omp_for_combined_into_p (fd->for_stmt))
     {
-      struct omp_context *task_ctx = NULL;
+      struct omp_context *taskreg_ctx = NULL;
       if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
 	{
 	  gomp_for *gfor = as_a <gomp_for *> (ctx->outer->stmt);
-	  if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR)
+	  if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR
+	      || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE)
 	    {
-	      struct omp_for_data outer_fd;
-	      extract_omp_for_data (gfor, &outer_fd, NULL);
-	      n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+	      if (gimple_omp_for_combined_into_p (gfor))
+		{
+		  gcc_assert (ctx->outer->outer
+			      && is_parallel_ctx (ctx->outer->outer));
+		  taskreg_ctx = ctx->outer->outer;
+		}
+	      else
+		{
+		  struct omp_for_data outer_fd;
+		  extract_omp_for_data (gfor, &outer_fd, NULL);
+		  n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+		}
 	    }
 	  else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP)
-	    task_ctx = ctx->outer->outer;
+	    taskreg_ctx = ctx->outer->outer;
 	}
-      else if (is_task_ctx (ctx->outer))
-	task_ctx = ctx->outer;
-      if (task_ctx)
+      else if (is_taskreg_ctx (ctx->outer))
+	taskreg_ctx = ctx->outer;
+      if (taskreg_ctx)
 	{
 	  int i;
 	  tree innerc
-	    = find_omp_clause (gimple_omp_task_clauses (task_ctx->stmt),
+	    = find_omp_clause (gimple_omp_taskreg_clauses (taskreg_ctx->stmt),
 			       OMP_CLAUSE__LOOPTEMP_);
 	  gcc_assert (innerc);
 	  for (i = 0; i < fd->collapse; i++)
@@ -13489,7 +13550,7 @@ lower_omp_for_lastprivate (struct omp_fo
 	  if (innerc)
 	    n2 = fold_convert (TREE_TYPE (n2),
 			       lookup_decl (OMP_CLAUSE_DECL (innerc),
-					    task_ctx));
+					    taskreg_ctx));
 	}
     }
   cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
--- gcc/c-family/c-omp.c.jj	2015-10-22 13:41:23.256881650 +0200
+++ gcc/c-family/c-omp.c	2015-10-22 14:35:41.075032703 +0200
@@ -1097,10 +1097,24 @@ c_omp_split_clauses (location_t loc, enu
 	      s = C_OMP_CLAUSE_SPLIT_FOR;
 	    }
 	  break;
-	/* Lastprivate is allowed on for, sections and simd.  In
+	/* Lastprivate is allowed on distribute, for, sections and simd.  In
 	   parallel {for{, simd},sections} we actually want to put it on
 	   parallel rather than for or sections.  */
 	case OMP_CLAUSE_LASTPRIVATE:
+	  if (code == OMP_DISTRIBUTE)
+	    {
+	      s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE;
+	      break;
+	    }
+	  if ((mask & (OMP_CLAUSE_MASK_1
+		       << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)) != 0)
+	    {
+	      c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+				    OMP_CLAUSE_LASTPRIVATE);
+	      OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+	      OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE];
+	      cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE] = c;
+	    }
 	  if (code == OMP_FOR || code == OMP_SECTIONS)
 	    {
 	      if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS))
--- gcc/c/c-parser.c.jj	2015-10-22 13:41:23.230882028 +0200
+++ gcc/c/c-parser.c	2015-10-22 14:35:41.075032703 +0200
@@ -14682,6 +14682,7 @@ c_parser_omp_cancellation_point (c_parse
 #define OMP_DISTRIBUTE_CLAUSE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
 
--- gcc/cp/parser.c.jj	2015-10-22 13:41:23.248881766 +0200
+++ gcc/cp/parser.c	2015-10-22 14:35:41.067032818 +0200
@@ -33572,6 +33572,7 @@ cp_parser_omp_cancellation_point (cp_par
 #define OMP_DISTRIBUTE_CLAUSE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
 
--- gcc/testsuite/c-c++-common/gomp/distribute-1.c.jj	2015-10-22 14:35:41.068032804 +0200
+++ gcc/testsuite/c-c++-common/gomp/distribute-1.c	2015-10-22 14:35:41.068032804 +0200
@@ -0,0 +1,56 @@
+int s1, s2, s3, s4, s5, s6, s7, s8;
+#pragma omp declare target (s1, s2, s3, s4, s5, s6, s7, s8)
+
+void
+f1 (void)
+{
+  int i;
+  #pragma omp distribute
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute private (i)
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute
+  for (int j = 0; j < 64; j++)
+    ;
+  #pragma omp distribute lastprivate (s1)
+  for (s1 = 0; s1 < 64; s1 += 2)
+    ;
+  #pragma omp distribute lastprivate (s2)
+  for (i = 0; i < 64; i++)
+    s2 = 2 * i;
+  #pragma omp distribute simd
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute simd lastprivate (s3, s4) collapse(2)
+  for (s3 = 0; s3 < 64; s3++)
+    for (s4 = 0; s4 < 3; s4++)
+      ;
+  #pragma omp distribute parallel for
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute parallel for private (i)
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute parallel for lastprivate (s5)
+  for (s5 = 0; s5 < 64; s5++)
+    ;
+  #pragma omp distribute firstprivate (s7) private (s8)
+  for (i = 0; i < 64; i++)
+    s8 = s7++;
+}
+
+void
+f2 (void)
+{
+  int i;
+  #pragma omp distribute lastprivate (i)	/* { dg-error "lastprivate variable .i. is private in outer context" } */
+  for (i = 0; i < 64; i++)
+    ;
+  #pragma omp distribute firstprivate (s6) lastprivate (s6) /* { dg-error "same variable used in .firstprivate. and .lastprivate. clauses on .distribute. construct" } */
+  for (i = 0; i < 64; i++)
+    s6 += i;
+}
+
+#pragma omp declare target to(f1, f2)
--- gcc/testsuite/c-c++-common/gomp/pr61486-2.c.jj	2015-10-22 13:41:22.983885625 +0200
+++ gcc/testsuite/c-c++-common/gomp/pr61486-2.c	2015-10-22 14:35:41.068032804 +0200
@@ -355,8 +355,11 @@ test (int n, int o, int p, int q, int r,
 
 int q, i, j;
 
+#pragma omp declare target
+int s;
+
 void
-test2 (int n, int o, int p, int r, int s, int *pp)
+test2 (int n, int o, int p, int r, int *pp)
 {
   int a[o];
     #pragma omp distribute collapse (2) dist_schedule (static, 4) firstprivate (q)
@@ -449,3 +452,4 @@ test2 (int n, int o, int p, int r, int s
 	  s = i * 10;
 	}
 }
+#pragma omp end declare target
--- libgomp/testsuite/libgomp.c/pr66199-5.c.jj	2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-5.c	2015-10-22 16:15:11.064658921 +0200
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b) shared(u, v, w)
+  for (d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+  return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-6.c.jj	2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-6.c	2015-10-22 16:24:37.000000000 +0200
@@ -0,0 +1,42 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp" } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams distribute parallel for default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-7.c.jj	2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-7.c	2015-10-22 17:11:33.419476232 +0200
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams distribute simd default(none) firstprivate (a, b) shared(u, v, w)
+  for (d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+  return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams distribute simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-8.c.jj	2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-8.c	2015-10-22 17:11:33.419476232 +0200
@@ -0,0 +1,70 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+  long d;
+  #pragma omp target map(from: d)
+  #pragma omp teams default(none) shared(a, b, d, u, v, w)
+  #pragma omp distribute simd firstprivate (a, b)
+  for (d = a; d < b; d++)
+    u[d] = v[d] + w[d];
+  return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  #pragma omp distribute simd linear(d) lastprivate(e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  #pragma omp distribute simd firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) firstprivate (a1, b1, a2, b2) shared(d1, d2, u, v, w)
+  #pragma omp distribute simd collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f1 (0, 1024) != 1024
+      || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64
+      || f4 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-9.c.jj	2015-10-22 17:11:56.362148829 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-9.c	2015-10-22 17:12:59.597246434 +0200
@@ -0,0 +1,43 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+  long d, e;
+  #pragma omp target map(from: d, e)
+  #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  #pragma omp distribute lastprivate(d, e)
+  for (d = a; d < b; d++)
+    {
+      u[d] = v[d] + w[d];
+      e = c + d * 5;
+    }
+  return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+  long d1, d2;
+  #pragma omp target map(from: d1, d2)
+  #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  #pragma omp distribute firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+  for (d1 = a1; d1 < b1; d1++)
+    for (d2 = a2; d2 < b2; d2++)
+      u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+  return d1 + d2;
+}
+
+int
+main ()
+{
+  if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+      || f3 (0, 32, 0, 32) != 64)
+    __builtin_abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c++/pr66199-3.C.jj	2015-10-22 17:14:28.277980917 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-3.C	2015-10-22 17:14:28.276980932 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-3.c"
--- libgomp/testsuite/libgomp.c++/pr66199-4.C.jj	2015-10-22 17:14:28.280980875 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-4.C	2015-10-22 17:14:28.279980889 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-4.c"
--- libgomp/testsuite/libgomp.c++/pr66199-5.C.jj	2015-10-22 17:14:28.283980832 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-5.C	2015-10-22 17:14:28.282980846 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-5.c"
--- libgomp/testsuite/libgomp.c++/pr66199-6.C.jj	2015-10-22 17:14:28.287980775 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-6.C	2015-10-22 17:14:28.285980803 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-6.c"
--- libgomp/testsuite/libgomp.c++/pr66199-7.C.jj	2015-10-22 17:14:28.290980732 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-7.C	2015-10-22 17:14:28.289980746 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-7.c"
--- libgomp/testsuite/libgomp.c++/pr66199-8.C.jj	2015-10-22 17:14:28.293980689 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-8.C	2015-10-22 17:14:28.292980703 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-8.c"
--- libgomp/testsuite/libgomp.c++/pr66199-9.C.jj	2015-10-22 17:14:28.296980646 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-9.C	2015-10-22 17:14:28.295980661 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-9.c"
