public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [gomp4] OpenACC first private
@ 2015-08-03 14:30 Nathan Sidwell
  2015-08-10 19:50 ` Thomas Schwinge
                   ` (2 more replies)
  0 siblings, 3 replies; 8+ messages in thread
From: Nathan Sidwell @ 2015-08-03 14:30 UTC (permalink / raw)
  To: GCC Patches; +Cc: james norris

[-- Attachment #1: Type: text/plain, Size: 498 bytes --]

I've committed this patch to gomp4.  The existing implementation of firstprivate 
presumes the existence of memory at the CTA level.  This patch does away with 
that, treating firstprivate as thread-private variables initialized from the 
host.

During development there was some fallout from declare handling, as that wasn't 
  creating the expected omp_region context object.  The previous handling of 
firstprivate just happened to work.  Jim has been working on resolving that problem.

nathan

[-- Attachment #2: gomp4-fp.patch --]
[-- Type: text/x-patch, Size: 17707 bytes --]

2015-08-03  Nathan Sidwell  <nathan@codesourcery.com>

	* gimplify.c (GOVD_GANGLOCAL): Delete.
	(oacc_default_clause): Only derereference reference types. Mark
	firstprivate as GOVD_FIRSTPRIVATE.
	(gimplify_adjust_omp_clauses_1): Remove GANGLOCALL handling.
	(gimplify_omp_for): Remove bogus OpenACC outer context lookup.
	* omp-low.c (build_outer_var_ref): Simplify openacc outer ref
	lookup.
	(scan_sharing_clauses): Handle openacc firstprivate.
	(lower_omp_target): Handle openacc firstprivate.

	c/
	* c-parser.c (c_parser_oacc_data_clause): Remove firstprivate
	handling.
	(c_parser_oac_all_clauses): Firstpribsste is a firstprivate
	clause.
	* c-typeck.c (c_finish_omp_clauses): Remove GANGLOCAL handling.

	fortran/
	* trans-openmp.c (gfc_trans_omp_clauses_1): Remove GANGLOCAL
	handling.
	* gfortran.h (OMP_MAP_GANGLOCAL): Delete.
	(OMP_MAP_FORCE_TO_GANGLOCAL):  Likewise.
	* openmp.c (gfc_match_omp_clauses): Remove openacc specific
	firstprivate handling.

	testsuite/
	* gfortran.dg/goacc/parallel-tree.f95: Remove ganglocal
	expectation.
	* gfortran.dg/goacc/list.f95: Stop expected firstprivate to be a
	data clause.
	* c-c++-common/goacc/firstprivate.c: Likewise.

	cp/
	* semantics.c (finish_omp_clauses): Remove OpenACC-specific
	firstprivate handling.
	* parser.c (cp_parser_oacc_data_clause): Remove firstprivate here.
	(cp_parser_oacc_all_clauses): First private is a firstprivate clause.

Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c	(revision 226462)
+++ gcc/gimplify.c	(working copy)
@@ -94,9 +94,6 @@ enum gimplify_omp_var_data
 
   GOVD_FORCE_MAP = 1 << 16,
 
-  /* Gang-local OpenACC variable.  */
-  GOVD_GANGLOCAL = 1 << 17,
-
   /* OpenACC deviceptr clause.  */
   GOVD_USE_DEVPTR = 1 << 18,
 
@@ -5937,14 +5934,13 @@ oacc_default_clause (struct gimplify_omp
 	if (is_global_var (decl) && device_resident_p (decl))
 	  flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
 	else if (ctx->acc_region_kind == ARK_KERNELS)
-	  /* Scalars under kernels are default 'copy'.  */
+	  /* Everything under kernels are default 'copy'.  */
 	  flags |= GOVD_FORCE_MAP | GOVD_MAP;
 	else if (ctx->acc_region_kind == ARK_PARALLEL)
 	  {
 	    tree type = TREE_TYPE (decl);
 
-	    /*  Should this  be REFERENCE_TYPE_P? */
-	    if (POINTER_TYPE_P (type))
+	    if (TREE_CODE (type) == REFERENCE_TYPE)
 	      type = TREE_TYPE (type);
 	
 	    if (AGGREGATE_TYPE_P (type))
@@ -5952,12 +5948,12 @@ oacc_default_clause (struct gimplify_omp
 	      flags |= GOVD_MAP;
 	    else
 	      /* Scalars default to 'firstprivate'.  */
-	      flags |= GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY | GOVD_MAP;
+	      flags |= GOVD_FIRSTPRIVATE;
 	  }
 	else
 	  gcc_unreachable ();
       }
-      break;
+    break;
     }
   
   return flags;
@@ -6812,10 +6808,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
   else if (code == OMP_CLAUSE_MAP)
     {
       OMP_CLAUSE_SET_MAP_KIND (clause,
-			       flags & GOVD_MAP_TO_ONLY
-			       ? (flags & GOVD_GANGLOCAL
-				  ? GOMP_MAP_FORCE_TO_GANGLOCAL
-				  : GOMP_MAP_TO)
+			       flags & GOVD_MAP_TO_ONLY ? GOMP_MAP_TO
 			       : (flags & GOVD_FORCE_MAP
 				  ? GOMP_MAP_FORCE_TOFROM
 				  : GOMP_MAP_TOFROM));
@@ -7542,11 +7535,7 @@ gimplify_omp_for (tree *expr_p, gimple_s
       else if (omp_is_private (gimplify_omp_ctxp, decl, 0))
 	omp_notice_variable (gimplify_omp_ctxp, decl, true);
       else
-	{
-	  if (ork == ORK_OACC && gimplify_omp_ctxp->outer_context)
-	    omp_notice_variable (gimplify_omp_ctxp->outer_context, decl, true);
-	  omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
-	}
+	omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
 
       /* If DECL is not a gimple register, create a temporary variable to act
 	 as an iteration counter.  This is valid, since DECL cannot be
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c	(revision 226462)
+++ gcc/c/c-parser.c	(working copy)
@@ -10719,9 +10719,6 @@ c_parser_oacc_data_clause (c_parser *par
     case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
       kind = GOMP_MAP_DEVICE_RESIDENT;
       break;
-    case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
-      kind = GOMP_MAP_FORCE_TO_GANGLOCAL;
-      break;
     case PRAGMA_OACC_CLAUSE_HOST:
       kind = GOMP_MAP_FORCE_FROM;
       break;
@@ -12316,7 +12313,7 @@ c_parser_oacc_all_clauses (c_parser *par
 	  c_name = "deviceptr";
 	  break;
 	case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
-	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  clauses = c_parser_omp_clause_firstprivate (parser, clauses);
 	  c_name = "firstprivate";
 	  break;
 	case PRAGMA_OACC_CLAUSE_GANG:
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c	(revision 226462)
+++ gcc/c/c-typeck.c	(working copy)
@@ -12435,10 +12435,6 @@ c_finish_omp_clauses (tree clauses, bool
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_TO_GANGLOCAL))
-		error_at (OMP_CLAUSE_LOCATION (c),
-			  "subarrays are not permitted in firstprivate");
 	      if (handle_omp_array_sections (c))
 		remove = true;
 	      else
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 226462)
+++ gcc/omp-low.c	(working copy)
@@ -1172,14 +1172,12 @@ build_outer_var_ref (tree var, omp_conte
       if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
 	  && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
 	{
-	  for (ctx = ctx->outer; ctx && !maybe_lookup_decl (var, ctx);
-	       ctx = ctx->outer)
-	    ;
-
-	  if (ctx == NULL)
-	    gcc_unreachable ();
-
-	  x = lookup_decl (var, ctx);
+	  do
+	    {
+	      ctx = ctx->outer;
+	      x = maybe_lookup_decl (var, ctx);
+	    }
+	  while(!x);
 	}
       else
 	x = lookup_decl (var, ctx->outer);
@@ -1848,10 +1846,6 @@ scan_sharing_clauses (tree clauses, omp_
 	  /* FALLTHRU */
 
 	case OMP_CLAUSE_FIRSTPRIVATE:
-	  if (is_gimple_omp_oacc (ctx->stmt))
-	    /* Clause represented by a gang-local map under OpenACC.  */
-	    gcc_unreachable ();
-	  /* FALLTHRU */
 	case OMP_CLAUSE_REDUCTION:
 	case OMP_CLAUSE_LINEAR:
 	  decl = OMP_CLAUSE_DECL (c);
@@ -1879,10 +1873,20 @@ scan_sharing_clauses (tree clauses, omp_
 	      else if (!global)
 		install_var_field (decl, by_ref, 3, ctx);
 	    }
-	  /* The gimplifier always includes a OMP_CLAUSE_MAP with each parallel
-	     reduction variable.  So don't install a local variable here.  */
+
 	  if (!is_oacc_parallel (ctx))
 	    install_var_local (decl, ctx);
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+	    {
+	      install_var_field (decl, (TREE_CODE (TREE_TYPE (decl))
+					!= REFERENCE_TYPE), 3, ctx);
+	      install_var_local (decl, ctx);
+	    }
+	  else
+	    /* The gimplifier always includes a OMP_CLAUSE_MAP with
+	       each parallel reduction variable.  So don't install a
+	       local variable here.  */
+	    gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION);
 	  break;
 
 	case OMP_CLAUSE__LOOPTEMP_:
@@ -2063,12 +2067,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_REDUCTION:
 	case OMP_CLAUSE_LINEAR:
@@ -11712,7 +11710,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;
-  gimple_seq tgt_body, olist, ilist, orlist, irlist, new_body;
+  gimple_seq tgt_body, olist, ilist, orlist, irlist, fplist, new_body;
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region, has_reduction;
   unsigned int map_cnt = 0;
@@ -11764,6 +11762,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))
@@ -11772,6 +11771,11 @@ lower_omp_target (gimple_stmt_iterator *
 
       default:
 	break;
+      case OMP_CLAUSE_FIRSTPRIVATE:
+	if (is_oacc_parallel (ctx))
+	  goto first_private;
+	break;
+	
       case OMP_CLAUSE_MAP:
 #ifdef ENABLE_CHECKING
 	/* First check what we're prepared to handle in the following.  */
@@ -11803,6 +11807,8 @@ lower_omp_target (gimple_stmt_iterator *
 	  /* FALLTHRU */
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
+      first_private:
+	
 	var = OMP_CLAUSE_DECL (c);
 	if (!DECL_P (var))
 	  {
@@ -11829,11 +11835,26 @@ 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);
-	    if (DECL_P (new_var))
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+	      {
+		if (TREE_CODE (TREE_TYPE (new_var)) == REFERENCE_TYPE)
+		  {
+		    /* 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;
@@ -11856,6 +11877,7 @@ lower_omp_target (gimple_stmt_iterator *
 	      }
 	  }
 	map_cnt++;
+	break;
       }
 
   if (offloaded)
@@ -11945,6 +11967,10 @@ lower_omp_target (gimple_stmt_iterator *
 
 	  default:
 	    break;
+	  case OMP_CLAUSE_FIRSTPRIVATE:
+	    if (!is_oacc_parallel (ctx))
+	      break;
+	    /* FALLTHROUGH */
 	  case OMP_CLAUSE_MAP:
 	  case OMP_CLAUSE_TO:
 	  case OMP_CLAUSE_FROM:
@@ -12011,6 +12037,14 @@ 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)
+		  {
+		    if (TREE_CODE (TREE_TYPE (var)) != REFERENCE_TYPE)
+		      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);
@@ -12039,7 +12073,16 @@ lower_omp_target (gimple_stmt_iterator *
 		    gimplify_assign (x, var, &ilist);
 		  }
 	      }
-	    tree s = OMP_CLAUSE_SIZE (c);
+	    tree s = NULL_TREE;
+	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FIRSTPRIVATE)
+	      s = OMP_CLAUSE_SIZE (c);
+	    else
+	      {
+		s = TREE_TYPE (ovar);
+		if (TREE_CODE (s) == REFERENCE_TYPE)
+		  s = TREE_TYPE (s);
+		s = TYPE_SIZE_UNIT (s);
+	      }
 	    if (s == NULL_TREE)
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
 	    s = fold_convert (size_type_node, s);
@@ -12054,6 +12097,9 @@ lower_omp_target (gimple_stmt_iterator *
 	      case OMP_CLAUSE_MAP:
 		tkind = OMP_CLAUSE_MAP_KIND (c);
 		break;
+	      case OMP_CLAUSE_FIRSTPRIVATE:
+		tkind = GOMP_MAP_TO;
+		break;
 	      case OMP_CLAUSE_TO:
 		tkind = GOMP_MAP_TO;
 		break;
@@ -12118,6 +12164,7 @@ lower_omp_target (gimple_stmt_iterator *
 	  		   gimple_build_assign (ctx->receiver_decl, t));
     }
   gimple_seq_add_seq (&new_body, ctx->ganglocal_init);
+  gimple_seq_add_seq (&new_body, fplist);
 
   if (offloaded)
     {
Index: gcc/fortran/trans-openmp.c
===================================================================
--- gcc/fortran/trans-openmp.c	(revision 226462)
+++ gcc/fortran/trans-openmp.c	(working copy)
@@ -2125,9 +2125,6 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl
 		case OMP_MAP_FROM:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
 		  break;
-		case OMP_MAP_GANGLOCAL:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
-		  break;
 		case OMP_MAP_TOFROM:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
 		  break;
@@ -2152,9 +2149,6 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl
 		case OMP_MAP_FORCE_DEVICEPTR:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
 		  break;
-		case OMP_MAP_FORCE_TO_GANGLOCAL:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
-		  break;
 		case OMP_MAP_DEVICE_RESIDENT:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DEVICE_RESIDENT);
 		  break;
Index: gcc/fortran/gfortran.h
===================================================================
--- gcc/fortran/gfortran.h	(revision 226462)
+++ gcc/fortran/gfortran.h	(working copy)
@@ -1138,7 +1138,6 @@ typedef enum
   OMP_MAP_ALLOC,
   OMP_MAP_TO,
   OMP_MAP_FROM,
-  OMP_MAP_GANGLOCAL,
   OMP_MAP_TOFROM,
   OMP_MAP_FORCE_ALLOC,
   OMP_MAP_FORCE_DEALLOC,
@@ -1149,7 +1148,6 @@ typedef enum
   OMP_MAP_FORCE_DEVICEPTR,
   OMP_MAP_DEVICE_RESIDENT,
   OMP_MAP_LINK,
-  OMP_MAP_FORCE_TO_GANGLOCAL
 }
 gfc_omp_map_op;
 
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c	(revision 226462)
+++ gcc/fortran/openmp.c	(working copy)
@@ -586,22 +586,12 @@ gfc_match_omp_clauses (gfc_omp_clauses *
 					  &c->lists[OMP_LIST_PRIVATE], true)
 	     == MATCH_YES)
 	continue;
-      if (mask & OMP_CLAUSE_FIRSTPRIVATE)
-	{
-	  if (openacc)
-	    {
-	      if (gfc_match ("firstprivate ( ") == MATCH_YES
-		  && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					       OMP_MAP_GANGLOCAL, false))
-		continue;
-	    }
-	  else if (gfc_match_omp_variable_list ("firstprivate (",
+      if ((mask & OMP_CLAUSE_FIRSTPRIVATE)
+	  && gfc_match_omp_variable_list ("firstprivate (",
 					  &c->lists[OMP_LIST_FIRSTPRIVATE],
-						true)
-		   == MATCH_YES)
-	    continue;
-
-	}
+					  true)
+	      == MATCH_YES)
+	continue;
       if ((mask & OMP_CLAUSE_LASTPRIVATE)
 	  && gfc_match_omp_variable_list ("lastprivate (",
 					  &c->lists[OMP_LIST_LASTPRIVATE],
Index: gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95	(revision 226462)
+++ gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95	(working copy)
@@ -37,4 +37,3 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } 
-! { dg-final { scan-tree-dump-times "map\\(force_to_ganglocal:w" 1 "original" } }
Index: gcc/testsuite/gfortran.dg/goacc/list.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/list.f95	(revision 226462)
+++ gcc/testsuite/gfortran.dg/goacc/list.f95	(working copy)
@@ -5,7 +5,7 @@ program test
   implicit none
 
   integer :: i, j, k, l, a(10)
-  common /b/ j, k
+  common /b/ k
   real, pointer :: p1 => NULL()
   complex :: c, d(10)
 
@@ -64,8 +64,8 @@ program test
 
   !$acc parallel firstprivate(10) ! { dg-error "Syntax error" }
 
-  !$acc parallel firstprivate (/b/, /b/) ! { dg-error "Syntax error" }
-  !$acc end parallel ! { dg-error "Unexpected" }
+  !$acc parallel firstprivate (/b/, /b/) ! { dg-error "present on multiple clauses" }
+  !$acc end parallel
 
   !$acc parallel firstprivate (i, j, i) ! { dg-error "present on multiple clauses" }
   !$acc end parallel
Index: gcc/testsuite/c-c++-common/goacc/firstprivate.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/firstprivate.c	(revision 226462)
+++ gcc/testsuite/c-c++-common/goacc/firstprivate.c	(working copy)
@@ -4,6 +4,6 @@ foo (void)
   int a, b[100];
 #pragma acc parallel firstprivate (a, b)
     ;
-#pragma acc parallel firstprivate (b[10:20]) /* { dg-error "subarrays are not permitted in firstprivate" } */
+#pragma acc parallel firstprivate (b[10:20]) /* { dg-error "expected" } */
     ;
 }
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c	(revision 226462)
+++ gcc/cp/semantics.c	(working copy)
@@ -5838,10 +5838,6 @@ finish_omp_clauses (tree clauses, bool o
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_TO_GANGLOCAL))
-		error_at (OMP_CLAUSE_LOCATION (c),
-			  "subarrays are not permitted in firstprivate");
 	      if (handle_omp_array_sections (c))
 		remove = true;
 	      else
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c	(revision 226462)
+++ gcc/cp/parser.c	(working copy)
@@ -28195,9 +28195,6 @@ cp_parser_oacc_data_clause (cp_parser *p
     case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
       kind = GOMP_MAP_DEVICE_RESIDENT;
       break;
-    case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
-      kind = GOMP_MAP_FORCE_TO_GANGLOCAL;
-      break;
     case PRAGMA_OACC_CLAUSE_HOST:
       kind = GOMP_MAP_FORCE_FROM;
       break;
@@ -29753,7 +29750,8 @@ cp_parser_oacc_all_clauses (cp_parser *p
 	  c_name = "deviceptr";
 	  break;
 	case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
-	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+	  clauses = cp_parser_omp_var_list
+	    (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses);
 	  c_name = "firstprivate";
 	  break;
 	case PRAGMA_OACC_CLAUSE_IF:

^ permalink raw reply	[flat|nested] 8+ messages in thread
end of thread, other threads:[~2015-10-29  8:23 UTC | newest]

Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-08-03 14:30 [gomp4] OpenACC first private Nathan Sidwell
2015-08-10 19:50 ` Thomas Schwinge
2015-08-18 23:30   ` Thomas Schwinge
2015-08-19 12:41     ` Nathan Sidwell
2015-08-12 18:31 ` [gomp4] declare directive James Norris
2015-09-01 16:34   ` Tom de Vries
2015-09-02  8:25   ` Tom de Vries
2015-10-29  8:31 ` [gomp4] OpenACC first private Thomas Schwinge

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).