diff mbox

[gomp4] OpenACC first private

Message ID 55BF7B19.7000409@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Aug. 3, 2015, 2:30 p.m. UTC
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
diff mbox

Patch

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: