diff mbox

OpenACC Firstprivate

Message ID 5641FB67.9010409@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Nov. 10, 2015, 2:12 p.m. UTC
On 11/09/15 09:46, Nathan Sidwell wrote:

> I'm going to try and get clarification, but I think the intent is to initialize
> with the value seen on the device.  Consider:

> My thinking is that the intent of the firstprivate is to initialize with the
> value known on the device (and behave as-if copyin, if it's not there).  Not the
> value most recently seen on the host -- the update clause could change that, and
> may well be being used as a debugging aide, so it seems bizarre that it can
> change program semantics in such a way.

We believe my example is well formed.  The data clauses transfer liveness of the 
data from host to device (and vice versa).  It is ill formed to manipulate the 
data on the non-live system.  firstprivate's intial value is taken from the 
(statically determined) live location.

Unless I'm misunderstanding something about GOMP_MAP_FIRSTPRIVATE, using regular 
target mapping is the right thing.

Here's an updated patch with the other two issues you noted fixed.


nathan

Comments

Jakub Jelinek Nov. 11, 2015, 8:04 a.m. UTC | #1
On Tue, Nov 10, 2015 at 09:12:55AM -0500, Nathan Sidwell wrote:
> +		    /* Create a local object to hold the instance
> +		       value.  */
> +		    tree inst = create_tmp_var
> +		      (TREE_TYPE (TREE_TYPE (new_var)),
> +		       IDENTIFIER_POINTER (DECL_NAME (new_var)));

Can you please rewrite this as:
		    tree type = TREE_TYPE (TREE_TYPE (new_var));
		    tree n = DECL_NAME (new_var);
		    tree inst = create_tmp_var (type, IDENTIFIER_POINTER (n));
or so (perhaps
		    const char *name
		      = IDENTIFIER_POINTER (DECL_NAME (new_var));
instead but then it takes one more line)?
I really don't like line breaks before opening ( unless really
necessary.

Otherwise LGTM.

	Jakub
Nathan Sidwell Nov. 11, 2015, 1:44 p.m. UTC | #2
On 11/11/15 03:04, Jakub Jelinek wrote:
> On Tue, Nov 10, 2015 at 09:12:55AM -0500, Nathan Sidwell wrote:
>> +		    /* Create a local object to hold the instance
>> +		       value.  */
>> +		    tree inst = create_tmp_var
>> +		      (TREE_TYPE (TREE_TYPE (new_var)),
>> +		       IDENTIFIER_POINTER (DECL_NAME (new_var)));
>
> Can you please rewrite this as:
> 		    tree type = TREE_TYPE (TREE_TYPE (new_var));
> 		    tree n = DECL_NAME (new_var);
> 		    tree inst = create_tmp_var (type, IDENTIFIER_POINTER (n));
> or so (perhaps
> 		    const char *name
> 		      = IDENTIFIER_POINTER (DECL_NAME (new_var));
> instead but then it takes one more line)?
> I really don't like line breaks before opening ( unless really
> necessary.

Oh, yeah you mentioned that before :)

>
> Otherwise LGTM.

thanks.

nathan
Thomas Schwinge Nov. 12, 2015, 8:59 a.m. UTC | #3
Hi Nathan!

Merging back your trunk r230169 into gomp-4_0-branch, for the new
libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c test, I'm
seeing the compiler diagnose as follows (compile with "-Wall -O2"):

    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: In function 'main._omp_fn.1':
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c:20:17: warning: 'val' is used uninitialized in this function [-Wuninitialized]
           ok  = val == 7;
                     ^
    
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c:9:7: note: 'val' was declared here
       int val = 2;
           ^

..., and execution fails ("return 1" from main), so I XFAILed the
execution in the merge commit r230214 on gomp-4_0-branch.  (..., and I
still think that it's a good idea to change the libgomp testsuite to run
with -Wall enabled...)

Do you have an idea what's going on?  Given your preparatory "[gomp4]
Rework gimplifyier region flags",
<http://news.gmane.org/find-root.php?message_id=%3C56434833.7010703%40acm.org%3E>
(thanks!), the merge commit r230214 on gomp-4_0-branch didn't contain any
changes to gcc/gimplify.c, so that can't be it.  It also can't be the
possibly inconsistent usage of gcc/omp-low.c:is_reference vs. "TREE_CODE
(TREE_TYPE ([...])) == REFERENCE_TYPE" in gcc/omp-low.c, because that
doesn't matter for C code anyway (no artificial REFERENCE_TYPEs
generated), right?  So it must be some other change installed on
gomp-4_0-branch but not on trunk.


Grüße
 Thomas
diff mbox

Patch

2015-11-10  Nathan Sidwell  <nathan@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* gcc/gimplify.c (enum  omp_region_type): Add ORT_ACC,
	ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS.  Adjust ORT_NONE.
	(gimple_add_tmp_var): Add ORT_ACC checks.
	(gimplify_var_or_parm_decl): Likewise.
	(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
	mask.
	(omp_add_variable): Look in outer contexts for openacc and allow
	reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
	checks.
	(omp_notice_variable, omp_is_private, omp_check_private): Add
	ORT_ACC checks.
	(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
	Permit private openacc reductions.
	(gimplify_oacc_cache): Specify ORT_ACC.
	(gimplify_omp_workshare): Adjust OpenACC region types.
	(gimplify_omp_target_update): Likewise.
	* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
	firstprivate sorry.
	(lower-rec_input_clauses): Don't handle openacc firstprivate
	references here.
	(lower_omp_target): Emit initializers for openacc firstprivate vars.

	gcc/testsuite/
	* gfortran.dg/goacc/private-3.f95: Remove xfail.
	* gfortran.dg/goacc/combined_loop.f90: Remove xfail.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.

Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c	(revision 230107)
+++ gcc/gimplify.c	(working copy)
@@ -95,22 +95,34 @@  enum gimplify_omp_var_data
 
 enum omp_region_type
 {
-  ORT_WORKSHARE = 0,
-  ORT_SIMD = 1,
-  ORT_PARALLEL = 2,
-  ORT_COMBINED_PARALLEL = 3,
-  ORT_TASK = 4,
-  ORT_UNTIED_TASK = 5,
-  ORT_TEAMS = 8,
-  ORT_COMBINED_TEAMS = 9,
+  ORT_WORKSHARE = 0x00,
+  ORT_SIMD 	= 0x01,
+
+  ORT_PARALLEL	= 0x02,
+  ORT_COMBINED_PARALLEL = 0x03,
+
+  ORT_TASK	= 0x04,
+  ORT_UNTIED_TASK = 0x05,
+
+  ORT_TEAMS	= 0x08,
+  ORT_COMBINED_TEAMS = 0x09,
+
   /* Data region.  */
-  ORT_TARGET_DATA = 16,
+  ORT_TARGET_DATA = 0x10,
+
   /* Data region with offloading.  */
-  ORT_TARGET = 32,
-  ORT_COMBINED_TARGET = 33,
+  ORT_TARGET	= 0x20,
+  ORT_COMBINED_TARGET = 0x21,
+
+  /* OpenACC variants.  */
+  ORT_ACC	= 0x40,  /* A generic OpenACC region.  */
+  ORT_ACC_DATA	= ORT_ACC | ORT_TARGET_DATA, /* Data construct.  */
+  ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET,  /* Parallel construct */
+  ORT_ACC_KERNELS  = ORT_ACC | ORT_TARGET | 0x80,  /* Kernels construct.  */
+
   /* Dummy OpenMP region, used to disable expansion of
      DECL_VALUE_EXPRs in taskloop pre body.  */
-  ORT_NONE = 64
+  ORT_NONE	= 0x100
 };
 
 /* Gimplify hashtable helper.  */
@@ -689,7 +701,8 @@  gimple_add_tmp_var (tree tmp)
 	  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
 	  while (ctx
 		 && (ctx->region_type == ORT_WORKSHARE
-		     || ctx->region_type == ORT_SIMD))
+		     || ctx->region_type == ORT_SIMD
+		     || ctx->region_type == ORT_ACC))
 	    ctx = ctx->outer_context;
 	  if (ctx)
 	    omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
@@ -1804,7 +1817,8 @@  gimplify_var_or_parm_decl (tree *expr_p)
 	  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
 	  while (ctx
 		 && (ctx->region_type == ORT_WORKSHARE
-		     || ctx->region_type == ORT_SIMD))
+		     || ctx->region_type == ORT_SIMD
+		     || ctx->region_type == ORT_ACC))
 	    ctx = ctx->outer_context;
 	  if (!ctx && !nonlocal_vlas->add (decl))
 	    {
@@ -5579,7 +5593,8 @@  omp_firstprivatize_variable (struct gimp
 	}
       else if (ctx->region_type != ORT_WORKSHARE
 	       && ctx->region_type != ORT_SIMD
-	       && ctx->region_type != ORT_TARGET_DATA)
+	       && ctx->region_type != ORT_ACC
+	       && !(ctx->region_type & ORT_TARGET_DATA))
 	omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
 
       ctx = ctx->outer_context;
@@ -5667,11 +5682,13 @@  omp_add_variable (struct gimplify_omp_ct
       /* We shouldn't be re-adding the decl with the same data
 	 sharing class.  */
       gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
-      /* The only combination of data sharing classes we should see is
-	 FIRSTPRIVATE and LASTPRIVATE.  */
       nflags = n->value | flags;
-      gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
-		  == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
+      /* The only combination of data sharing classes we should see is
+	 FIRSTPRIVATE and LASTPRIVATE.  However, OpenACC permits
+	 reduction variables to be used in data sharing clauses.  */
+      gcc_assert ((ctx->region_type & ORT_ACC) != 0
+		  || ((nflags & GOVD_DATA_SHARE_CLASS)
+		      == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
 		  || (flags & GOVD_DATA_SHARE_CLASS) == 0);
       n->value = nflags;
       return;
@@ -5968,20 +5985,47 @@  omp_notice_variable (struct gimplify_omp
 	      else if (is_scalar)
 		nflags |= GOVD_FIRSTPRIVATE;
 	    }
-	  tree type = TREE_TYPE (decl);
-	  if (nflags == flags
-	      && gimplify_omp_ctxp->target_firstprivatize_array_bases
-	      && lang_hooks.decls.omp_privatize_by_reference (decl))
-	    type = TREE_TYPE (type);
-	  if (nflags == flags
-	      && !lang_hooks.types.omp_mappable_type (type))
-	    {
-	      error ("%qD referenced in target region does not have "
-		     "a mappable type", decl);
-	      nflags |= GOVD_MAP | GOVD_EXPLICIT;
+
+	  struct gimplify_omp_ctx *octx = ctx->outer_context;
+	  if ((ctx->region_type & ORT_ACC) && octx)
+	    {
+	      /* Look in outer OpenACC contexts, to see if there's a
+		 data attribute for this variable.  */
+	      omp_notice_variable (octx, decl, in_code);
+
+	      for (; octx; octx = octx->outer_context)
+		{
+		  if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
+		    break;
+		  splay_tree_node n2
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  if (n2)
+		    {
+		      nflags |= GOVD_MAP;
+		      goto found_outer;
+		    }
+		}
 	    }
-	  else if (nflags == flags)
-	    nflags |= GOVD_MAP;
+
+	  {
+	    tree type = TREE_TYPE (decl);
+
+	    if (nflags == flags
+		&& gimplify_omp_ctxp->target_firstprivatize_array_bases
+		&& lang_hooks.decls.omp_privatize_by_reference (decl))
+	      type = TREE_TYPE (type);
+	    if (nflags == flags
+		&& !lang_hooks.types.omp_mappable_type (type))
+	      {
+		error ("%qD referenced in target region does not have "
+		       "a mappable type", decl);
+		nflags |= GOVD_MAP | GOVD_EXPLICIT;
+	      }
+	    else if (nflags == flags)
+	      nflags |= GOVD_MAP;
+	  }
+	found_outer:
 	  omp_add_variable (ctx, decl, nflags);
 	}
       else
@@ -5998,7 +6042,8 @@  omp_notice_variable (struct gimplify_omp
     {
       if (ctx->region_type == ORT_WORKSHARE
 	  || ctx->region_type == ORT_SIMD
-	  || ctx->region_type == ORT_TARGET_DATA)
+	  || ctx->region_type == ORT_ACC
+	  || (ctx->region_type & ORT_TARGET_DATA) != 0)
 	goto do_outer;
 
       flags = omp_default_clause (ctx, decl, in_code, flags);
@@ -6112,7 +6157,8 @@  omp_is_private (struct gimplify_omp_ctx
     }
 
   if (ctx->region_type != ORT_WORKSHARE
-      && ctx->region_type != ORT_SIMD)
+      && ctx->region_type != ORT_SIMD
+      && ctx->region_type != ORT_ACC)
     return false;
   else if (ctx->outer_context)
     return omp_is_private (ctx->outer_context, decl, simd);
@@ -6168,7 +6214,8 @@  omp_check_private (struct gimplify_omp_c
 	}
     }
   while (ctx->region_type == ORT_WORKSHARE
-	 || ctx->region_type == ORT_SIMD);
+	 || ctx->region_type == ORT_SIMD
+	 || ctx->region_type == ORT_ACC);
   return false;
 }
 
@@ -6311,7 +6358,8 @@  gimplify_scan_omp_clauses (tree *list_p,
 		omp_notice_variable (outer_ctx->outer_context, decl, true);
 	    }
 	  else if (outer_ctx
-		   && outer_ctx->region_type == ORT_WORKSHARE
+		   && (outer_ctx->region_type == ORT_WORKSHARE
+		       || outer_ctx->region_type == ORT_ACC)
 		   && outer_ctx->combined_loop
 		   && splay_tree_lookup (outer_ctx->variables,
 					 (splay_tree_key) decl) == NULL
@@ -6335,7 +6383,9 @@  gimplify_scan_omp_clauses (tree *list_p,
 	  goto do_add;
 	case OMP_CLAUSE_REDUCTION:
 	  flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
-	  check_non_private = "reduction";
+	  /* OpenACC permits reductions on private variables.  */
+	  if (!(region_type & ORT_ACC))
+	    check_non_private = "reduction";
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (decl) == MEM_REF)
 	    {
@@ -7704,7 +7754,7 @@  gimplify_oacc_cache (tree *expr_p, gimpl
 {
   tree expr = *expr_p;
 
-  gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
+  gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC,
 			     OACC_CACHE);
   gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
 
@@ -7833,7 +7883,9 @@  gimplify_omp_for (tree *expr_p, gimple_s
     case OMP_FOR:
     case CILK_FOR:
     case OMP_DISTRIBUTE:
+      break;
     case OACC_LOOP:
+      ort = ORT_ACC;
       break;
     case OMP_TASKLOOP:
       if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
@@ -8895,10 +8947,14 @@  gimplify_omp_workshare (tree *expr_p, gi
       ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
       break;
     case OACC_KERNELS:
+      ort = ORT_ACC_KERNELS;
+      break;
     case OACC_PARALLEL:
-      ort = ORT_TARGET;
+      ort = ORT_ACC_PARALLEL;
       break;
     case OACC_DATA:
+      ort = ORT_ACC_DATA;
+      break;
     case OMP_TARGET_DATA:
       ort = ORT_TARGET_DATA;
       break;
@@ -8920,7 +8976,7 @@  gimplify_omp_workshare (tree *expr_p, gi
 	pop_gimplify_context (g);
       else
 	pop_gimplify_context (NULL);
-      if (ort == ORT_TARGET_DATA)
+      if ((ort & ORT_TARGET_DATA) != 0)
 	{
 	  enum built_in_function end_ix;
 	  switch (TREE_CODE (expr))
@@ -8995,17 +9051,18 @@  gimplify_omp_target_update (tree *expr_p
   tree expr = *expr_p;
   int kind;
   gomp_target *stmt;
+  enum omp_region_type ort = ORT_WORKSHARE;
 
   switch (TREE_CODE (expr))
     {
     case OACC_ENTER_DATA:
-      kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
-      break;
     case OACC_EXIT_DATA:
       kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+      ort = ORT_ACC;
       break;
     case OACC_UPDATE:
       kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+      ort = ORT_ACC;
       break;
     case OMP_TARGET_UPDATE:
       kind = GF_OMP_TARGET_KIND_UPDATE;
@@ -9020,7 +9077,7 @@  gimplify_omp_target_update (tree *expr_p
       gcc_unreachable ();
     }
   gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
-			     ORT_WORKSHARE, TREE_CODE (expr));
+			     ort, TREE_CODE (expr));
   gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
 			       TREE_CODE (expr));
   stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 230107)
+++ gcc/omp-low.c	(working copy)
@@ -1896,12 +1896,6 @@  scan_sharing_clauses (tree clauses, omp_
 	  /* FALLTHRU */
 
 	case OMP_CLAUSE_FIRSTPRIVATE:
-	  if (is_gimple_omp_oacc (ctx->stmt))
-	    {
-	      sorry ("clause not supported yet");
-	      break;
-	    }
-	  /* FALLTHRU */
 	case OMP_CLAUSE_LINEAR:
 	  decl = OMP_CLAUSE_DECL (c);
 	do_private:
@@ -2167,12 +2161,6 @@  scan_sharing_clauses (tree clauses, omp_
 	  /* FALLTHRU */
 
 	case OMP_CLAUSE_FIRSTPRIVATE:
-	  if (is_gimple_omp_oacc (ctx->stmt))
-	    {
-	      sorry ("clause not supported yet");
-	      break;
-	    }
-	  /* FALLTHRU */
 	case OMP_CLAUSE_PRIVATE:
 	case OMP_CLAUSE_LINEAR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
@@ -4684,7 +4672,7 @@  lower_rec_input_clauses (tree clauses, g
 		  gimplify_assign (ptr, x, ilist);
 		}
 	    }
-	  else if (is_reference (var))
+	  else if (is_reference (var) && !is_oacc_parallel (ctx))
 	    {
 	      /* For references that are being privatized for Fortran,
 		 allocate new backing storage for the new pointer
@@ -14911,7 +14899,7 @@  lower_omp_target (gimple_stmt_iterator *
   tree child_fn, t, c;
   gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
   gbind *tgt_bind, *bind, *dep_bind = NULL;
-  gimple_seq tgt_body, olist, ilist, new_body;
+  gimple_seq tgt_body, olist, ilist, fplist, new_body;
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
@@ -14963,6 +14951,7 @@  lower_omp_target (gimple_stmt_iterator *
   child_fn = ctx->cb.dst_fn;
 
   push_gimplify_context ();
+  fplist = NULL;
 
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
@@ -15007,6 +14996,7 @@  lower_omp_target (gimple_stmt_iterator *
 	  /* FALLTHRU */
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
+      oacc_firstprivate:
 	var = OMP_CLAUSE_DECL (c);
 	if (!DECL_P (var))
 	  {
@@ -15029,6 +15019,7 @@  lower_omp_target (gimple_stmt_iterator *
 	  }
 
 	if (offloaded
+	    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 	  {
@@ -15057,17 +15048,40 @@  lower_omp_target (gimple_stmt_iterator *
 	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
 
-	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
 	      x = build_simple_mem_ref (x);
-	    SET_DECL_VALUE_EXPR (new_var, x);
-	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+	      {
+		gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+		if (is_reference (new_var))
+		  {
+		    /* Create a local object to hold the instance
+		       value.  */
+		    tree inst = create_tmp_var
+		      (TREE_TYPE (TREE_TYPE (new_var)),
+		       IDENTIFIER_POINTER (DECL_NAME (new_var)));
+		    gimplify_assign (inst, fold_indirect_ref (x), &fplist);
+		    x = build_fold_addr_expr (inst);
+		  }
+		gimplify_assign (new_var, x, &fplist);
+	      }
+	    else if (DECL_P (new_var))
+	      {
+		SET_DECL_VALUE_EXPR (new_var, x);
+		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	      }
+	    else
+	      gcc_unreachable ();
 	  }
 	map_cnt++;
 	break;
 
       case OMP_CLAUSE_FIRSTPRIVATE:
+	if (is_oacc_parallel (ctx))
+	  goto oacc_firstprivate;
 	map_cnt++;
 	var = OMP_CLAUSE_DECL (c);
 	if (!is_reference (var)
@@ -15092,6 +15106,8 @@  lower_omp_target (gimple_stmt_iterator *
 	break;
 
       case OMP_CLAUSE_PRIVATE:
+	if (is_gimple_omp_oacc (ctx->stmt))
+	  break;
 	var = OMP_CLAUSE_DECL (c);
 	if (is_variable_sized (var))
 	  {
@@ -15195,9 +15211,11 @@  lower_omp_target (gimple_stmt_iterator *
 
 	  default:
 	    break;
+
 	  case OMP_CLAUSE_MAP:
 	  case OMP_CLAUSE_TO:
 	  case OMP_CLAUSE_FROM:
+	  oacc_firstprivate_map:
 	    nc = c;
 	    ovar = OMP_CLAUSE_DECL (c);
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -15248,9 +15266,9 @@  lower_omp_target (gimple_stmt_iterator *
 		x = build_sender_ref (ovar, ctx);
 
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-			 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
-			 && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
-			 && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+		    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		    && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+		    && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
 		  {
 		    gcc_assert (offloaded);
 		    tree avar
@@ -15261,6 +15279,15 @@  lower_omp_target (gimple_stmt_iterator *
 		    avar = build_fold_addr_expr (avar);
 		    gimplify_assign (x, avar, &ilist);
 		  }
+		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+		  {
+		    gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+		    if (!is_reference (var))
+		      var = build_fold_addr_expr (var);
+		    else
+		      talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
+		    gimplify_assign (x, var, &ilist);
+		  }
 		else if (is_gimple_reg (var))
 		  {
 		    gcc_assert (offloaded);
@@ -15289,7 +15316,17 @@  lower_omp_target (gimple_stmt_iterator *
 		    gimplify_assign (x, var, &ilist);
 		  }
 	      }
-	    s = OMP_CLAUSE_SIZE (c);
+	    s = NULL_TREE;
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+	      {
+		gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+		s = TREE_TYPE (ovar);
+		if (TREE_CODE (s) == REFERENCE_TYPE)
+		  s = TREE_TYPE (s);
+		s = TYPE_SIZE_UNIT (s);
+	      }
+	    else
+	      s = OMP_CLAUSE_SIZE (c);
 	    if (s == NULL_TREE)
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
 	    s = fold_convert (size_type_node, s);
@@ -15330,6 +15367,11 @@  lower_omp_target (gimple_stmt_iterator *
 		      tkind_zero = tkind;
 		  }
 		break;
+	      case OMP_CLAUSE_FIRSTPRIVATE:
+		gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+		tkind = GOMP_MAP_TO;
+		tkind_zero = tkind;
+		break;
 	      case OMP_CLAUSE_TO:
 		tkind = GOMP_MAP_TO;
 		tkind_zero = tkind;
@@ -15369,6 +15411,8 @@  lower_omp_target (gimple_stmt_iterator *
 	    break;
 
 	  case OMP_CLAUSE_FIRSTPRIVATE:
+	    if (is_oacc_parallel (ctx))
+	      goto oacc_firstprivate_map;
 	    ovar = OMP_CLAUSE_DECL (c);
 	    if (is_reference (ovar))
 	      talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@@ -15543,6 +15587,7 @@  lower_omp_target (gimple_stmt_iterator *
       gimple_seq_add_stmt (&new_body,
 	  		   gimple_build_assign (ctx->receiver_decl, t));
     }
+  gimple_seq_add_seq (&new_body, fplist);
 
   if (offloaded || data_region)
     {
@@ -15554,6 +15599,8 @@  lower_omp_target (gimple_stmt_iterator *
 	  default:
 	    break;
 	  case OMP_CLAUSE_FIRSTPRIVATE:
+	    if (is_gimple_omp_oacc (ctx->stmt))
+	      break;
 	    var = OMP_CLAUSE_DECL (c);
 	    if (is_reference (var)
 		|| is_gimple_reg_type (TREE_TYPE (var)))
@@ -15639,6 +15686,8 @@  lower_omp_target (gimple_stmt_iterator *
 	      }
 	    break;
 	  case OMP_CLAUSE_PRIVATE:
+	    if (is_gimple_omp_oacc (ctx->stmt))
+	      break;
 	    var = OMP_CLAUSE_DECL (c);
 	    if (is_reference (var))
 	      {
@@ -15727,7 +15776,7 @@  lower_omp_target (gimple_stmt_iterator *
       /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
 	 so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
 	 are already handled.  */
-      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+      for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
 	    tree var;
Index: gcc/testsuite/gfortran.dg/goacc/private-3.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/private-3.f95	(revision 230107)
+++ gcc/testsuite/gfortran.dg/goacc/private-3.f95	(working copy)
@@ -1,6 +1,4 @@ 
 ! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
 
 ! test for private variables in a reduction clause
 
Index: gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90	(revision 230107)
+++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90	(working copy)
@@ -1,6 +1,4 @@ 
 ! { dg-do compile } 
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
 
 !
 ! PR fortran/64726
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c	(revision 230107)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c	(working copy)
@@ -1,7 +1,5 @@ 
 /* { dg-do run } */
 /* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-   { dg-xfail-if "TODO" { *-*-* } } */
 
 #include <stdio.h>
 
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c	(revision 230107)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c	(working copy)
@@ -1,7 +1,5 @@ 
 /* { dg-do run } */
 /* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-   { dg-xfail-if "TODO" { *-*-* } } */
 
 #include <stdio.h>
 
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c	(revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c	(working copy)
@@ -0,0 +1,41 @@ 
+/* { dg-do run } */
+
+#include  <openacc.h>
+
+int main ()
+{
+  int ok = 1;
+  int val = 2;
+  int ary[32];
+  int ondev = 0;
+
+  for (int i = 0; i < 32; i++)
+    ary[i] = ~0;
+  
+#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev)
+  {
+    ondev = acc_on_device (acc_device_not_host);
+#pragma acc loop gang(static:1)
+    for (unsigned i = 0; i < 32; i++)
+      {
+	if (val != 2)
+	  ok = 0;
+	val += i;
+	ary[i] = val;
+      }
+  }
+
+  if (ondev)
+    {
+      if (!ok)
+	return 1;
+      if (val != 2)
+	return 1;
+
+      for (int i = 0; i < 32; i++)
+	if (ary[i] != 2 + i)
+	  return 1;
+    }
+  
+  return 0;
+}