[gomp4] openacc default handling

[gomp4] openacc default handling

[Nathan Sidwell]

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

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

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

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

Re: [gomp4] openacc default handling

[Thomas Schwinge]

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

Hi!

On Wed, 29 Jul 2015 19:13:01 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> 	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.

Committed to gomp-4_0-branch in r226495:

commit 9446d5a23e74653407f079188e8be4bc9343e15e
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Mon Aug 3 11:06:38 2015 +0000

    Resolve bootstrap failure in oacc_default_clause
    
        [...]/source-gcc/gcc/gimplify.c: In function 'unsigned int oacc_default_clause(gimplify_omp_ctx*, tree, bool, unsigned int)':
        [...]/source-gcc/gcc/gimplify.c:5913:13: error: unused parameter 'in_code' [-Werror=unused-parameter]
                bool in_code, unsigned flags)
                     ^
    
    Fixup for r226373.
    
    	gcc/
    	* gimplify.c (oacc_default_clause): Remove in_code formal
    	parameter.  Adjust all users.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@226495 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp |    5 +++++
 gcc/gimplify.c     |    5 ++---
 2 files changed, 7 insertions(+), 3 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 3b907ea..a30f6a3 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-08-03  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gimplify.c (oacc_default_clause): Remove in_code formal
+	parameter.  Adjust all users.
+
 2015-08-02  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* omp-low.c (oacc_xform_on_device): New function.
diff --git gcc/gimplify.c gcc/gimplify.c
index 0f8011d..f74f842 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -5909,8 +5909,7 @@ device_resident_p (tree decl)
    but not declared in an enclosing clause.  */
 
 static unsigned
-oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
-		     bool in_code, unsigned flags)
+oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 {
   switch (ctx->default_kind)
     {
@@ -6036,7 +6035,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	    }
 
 	  if (is_oacc)
-	    flags = oacc_default_clause (ctx, decl, in_code, flags);
+	    flags = oacc_default_clause (ctx, decl, flags);
 	  else
 	    flags |= GOVD_MAP;
 


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]