diff mbox

[gomp4] openacc default handling

Message ID 55B95DFD.6090703@acm.org
State New
Headers show

Commit Message

Nathan Sidwell July 29, 2015, 11:13 p.m. UTC 
I've committed this to  gomp4.  When I broke out the oacc_default_clause 
function from omp_notice_variable, I was puzzled by the ordering of the lookups, 
but could quite figure out what was wrong.  further investigation and standard 
reading showed that for  openacc, we look in outer lexical scopes to find a data 
clause for the object, before applying the default behaviour.  This patch 
reorganizes things to  make that the case.

nathan
diff mbox

Patch

2015-07-29  Nathan Sidwell  <nathan@codesourcery.com>

	gcc/
	* gimplify.c (oacc_default_clause): Outer scope searching moved to
	omp_notice_variable.
	(omp_notice_variable): For OpenACC search enclosing scopes before
	applying default.

	gcc/testsuite/
	* c-c++-common/goacc/default-2.c: New.

Index: gcc/testsuite/c-c++-common/goacc/default-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/default-2.c	(revision 0)
+++ gcc/testsuite/c-c++-common/goacc/default-2.c	(revision 0)
@@ -0,0 +1,21 @@ 
+void Foo ()
+{
+  int ary[10];
+  
+#pragma acc parallel default(none) /* { dg-error "enclosing" } */
+  {
+    ary[0] = 5; /* { dg-error "not specified" }  */
+  }
+}
+
+void Baz ()
+{
+  int ary[10];
+#pragma acc data copy (ary)
+  {
+#pragma acc parallel default(none)
+    {
+      ary[0] = 5;
+    }
+  }
+}
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c	(revision 226334)
+++ gcc/gimplify.c	(working copy)
@@ -5934,30 +5934,10 @@  oacc_default_clause (struct gimplify_omp
 
     case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
       {
-	if (struct gimplify_omp_ctx *octx = ctx->outer_context)
-	  {
-	    omp_notice_variable (octx, decl, in_code);
-	
-	    for (; octx; octx = octx->outer_context)
-	      {
-		if (octx->region_type & ORT_HOST_DATA)
-		  continue;
-		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)
-		{
-		  flags |= GOVD_MAP;
-		  goto found_outer;
-		}
-	      }
-	  }
-
 	if (is_global_var (decl) && device_resident_p (decl))
 	  flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
-	/* Scalars under kernels are default 'copy'.  */
 	else if (ctx->acc_region_kind == ARK_KERNELS)
+	  /* Scalars under kernels are default 'copy'.  */
 	  flags |= GOVD_FORCE_MAP | GOVD_MAP;
 	else if (ctx->acc_region_kind == ARK_PARALLEL)
 	  {
@@ -5968,16 +5948,14 @@  oacc_default_clause (struct gimplify_omp
 	      type = TREE_TYPE (type);
 	
 	    if (AGGREGATE_TYPE_P (type))
-	      /* Aggregates default to 'copy'.  This should really
-		 include GOVD_FORCE_MAP.  */
+	      /* Aggregates default to 'present_or_copy'.  */
 	      flags |= GOVD_MAP;
 	    else
-	      /* Scalars default tp 'firstprivate'.  */
+	      /* Scalars default to 'firstprivate'.  */
 	      flags |= GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY | GOVD_MAP;
 	  }
 	else
 	  gcc_unreachable ();
-      found_outer:;
       }
       break;
     }
@@ -6020,21 +5998,49 @@  omp_notice_variable (struct gimplify_omp
   if (ctx->region_type == ORT_TARGET)
     {
       ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
-      if (n == NULL)
+      bool is_oacc = ctx->region_kind == ORK_OACC;
+
+      if (!n)
 	{
-	  bool is_oacc = ctx->region_kind == ORK_OACC;
+	  struct gimplify_omp_ctx *octx = ctx->outer_context;
 
-	  if (is_oacc)
-	    flags = oacc_default_clause (ctx, decl, in_code, flags);
-	  else
-	    flags |= GOVD_MAP;
+	  /*  OpenMP doesn't look in outer contexts to find an
+	      enclosing data clause.  */
+	  if (is_oacc && octx)
+	    {
+	      omp_notice_variable (octx, decl, in_code);
+	      
+	      for (; octx; octx = octx->outer_context)
+		{
+		  if (octx->region_type & ORT_HOST_DATA)
+		    continue;
+		  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)
+		    {
+		      flags |= GOVD_MAP;
+		      goto found_outer;
+		    }
+		}
+	    }
 
-	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl), is_oacc))
+	  if (!lang_hooks.types.omp_mappable_type
+	      (TREE_TYPE (decl), ctx->region_kind == ORK_OACC))
 	    {
 	      error ("%qD referenced in target region does not have "
 		     "a mappable type", decl);
 	      flags |= GOVD_EXPLICIT;
 	    }
+
+	  if (is_oacc)
+	    flags = oacc_default_clause (ctx, decl, in_code, flags);
+	  else
+	    flags |= GOVD_MAP;
+
+	found_outer:;
 	  omp_add_variable (ctx, decl, flags);
 	}
       else