Re: Nested OpenACC/OpenMP constructs (was: OpenACC GIMPLE_OACC_* -- or not?)

[Thomas Schwinge <thomas@codesourcery.com>]

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

Hi Jakub!

On Wed, 10 Dec 2014 11:02:24 +0100, I wrote:
>     OpenACC: Rework nested constructs checking.
>     
>     	gcc/
>     	* omp-low.c (scan_omp_target): Remove taskreg_nesting_level and
>     	target_nesting_level assertions.
>     	(check_omp_nesting_restrictions): Rework OpenACC constructs
>     	handling.  Update and extend the relevant test cases.

Regarding the check_omp_nesting_restrictions rework:

> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -2706,46 +2700,26 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx)
>    scan_omp (gimple_omp_body_ptr (stmt), ctx);
>  }
>  
> -/* Check OpenMP nesting restrictions.  */
> +/* Check nesting restrictions.  */
>  static bool
>  check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
>  {
> -  /* TODO: While the OpenACC specification does allow for certain kinds of
> -     nesting, we don't support many of these yet.  */
> -  if (is_gimple_omp (stmt)
> -      && is_gimple_omp_oacc_specifically (stmt))
> +  /* TODO: Some OpenACC/OpenMP nesting should be allowed.  */
> +
> +  /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
> +     inside an OpenACC CTX.  */
> +  if (!(is_gimple_omp (stmt)
> +	&& is_gimple_omp_oacc_specifically (stmt)))
>      {
> -      /* Regular handling of OpenACC loop constructs.  */
> -      if (gimple_code (stmt) == GIMPLE_OMP_FOR
> -	  && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
> -	goto cont;
> -      /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different
> -	 from an OpenACC data construct.  */
> -      for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
> -	if (is_gimple_omp (ctx_->stmt)
> -	    && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET
> -		 && (gimple_omp_target_kind (ctx_->stmt)
> -		     == GF_OMP_TARGET_KIND_OACC_DATA)))
> -	  {
> -	    error_at (gimple_location (stmt),
> -		      "may not be nested");
> -	    return false;
> -	  }
> -    }
> -  else
> -    {
> -      /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP
> -	 builtin) inside any OpenACC CTX.  */
>        for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
>  	if (is_gimple_omp (ctx_->stmt)
>  	    && is_gimple_omp_oacc_specifically (ctx_->stmt))
>  	  {
>  	    error_at (gimple_location (stmt),
> -		      "may not be nested");
> +		      "non-OpenACC construct inside of OpenACC region");
>  	    return false;
>  	  }
>      }
> - cont:
>  
>    if (ctx != NULL)
>      {
> @@ -3003,20 +2977,74 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
>        break;
>      case GIMPLE_OMP_TARGET:
>        for (; ctx != NULL; ctx = ctx->outer)
> -	if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
> -	    && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
> -	  {
> -	    const char *name;
> -	    switch (gimple_omp_target_kind (stmt))
> -	      {
> -	      case GF_OMP_TARGET_KIND_REGION: name = "target"; break;
> -	      case GF_OMP_TARGET_KIND_DATA: name = "target data"; break;
> -	      case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break;
> -	      default: gcc_unreachable ();
> -	      }
> -	    warning_at (gimple_location (stmt), 0,
> -			"%s construct inside of target region", name);
> -	  }
> +	{
> +	  if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
> +	    {
> +	      if (is_gimple_omp (stmt)
> +		  && is_gimple_omp_oacc_specifically (stmt)
> +		  && is_gimple_omp (ctx->stmt))
> +		{
> +		  error_at (gimple_location (stmt),
> +			    "OpenACC construct inside of non-OpenACC region");
> +		  return false;
> +		}
> +	      continue;
> +	    }
> +
> +	  const char *stmt_name, *ctx_stmt_name;
> +	  switch (gimple_omp_target_kind (stmt))
> +	    {
> +	    case GF_OMP_TARGET_KIND_REGION: stmt_name = "target"; break;
> +	    case GF_OMP_TARGET_KIND_DATA: stmt_name = "target data"; break;
> +	    case GF_OMP_TARGET_KIND_UPDATE: stmt_name = "target update"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: stmt_name = "enter/exit data"; break;
> +	    default: gcc_unreachable ();
> +	    }
> +	  switch (gimple_omp_target_kind (ctx->stmt))
> +	    {
> +	    case GF_OMP_TARGET_KIND_REGION: ctx_stmt_name = "target"; break;
> +	    case GF_OMP_TARGET_KIND_DATA: ctx_stmt_name = "target data"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_PARALLEL: ctx_stmt_name = "parallel"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break;
> +	    case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
> +	    default: gcc_unreachable ();
> +	    }
> +
> +	  /* OpenACC/OpenMP mismatch?  */
> +	  if (is_gimple_omp_oacc_specifically (stmt)
> +	      != is_gimple_omp_oacc_specifically (ctx->stmt))
> +	    {
> +	      error_at (gimple_location (stmt),
> +			"%s %s construct inside of %s %s region",
> +			(is_gimple_omp_oacc_specifically (stmt)
> +			 ? "OpenACC" : "OpenMP"), stmt_name,
> +			(is_gimple_omp_oacc_specifically (ctx->stmt)
> +			 ? "OpenACC" : "OpenMP"), ctx_stmt_name);
> +	      return false;
> +	    }
> +	  if (is_gimple_omp_offloaded (ctx->stmt))
> +	    {
> +	      /* No GIMPLE_OMP_TARGET inside offloaded OpenACC CTX.  */
> +	      if (is_gimple_omp_oacc_specifically (ctx->stmt))
> +		{
> +		  error_at (gimple_location (stmt),
> +			    "%s construct inside of %s region",
> +			    stmt_name, ctx_stmt_name);
> +		  return false;
> +		}
> +	      else
> +		{
> +		  gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
> +		  warning_at (gimple_location (stmt), 0,
> +			      "%s construct inside of %s region",
> +			      stmt_name, ctx_stmt_name);
> +		}
> +	    }
> +	}
>        break;
>      default:
>        break;

For guarding against OpenMP target regressions, I used the following
tests -- OK to commit to trunk?

 gcc/testsuite/c-c++-common/gomp/nesting-1.c      | 77 ++++++++++++++++++++++++
 gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c | 23 +++++++
 2 files changed, 100 insertions(+)

diff --git gcc/testsuite/c-c++-common/gomp/nesting-1.c gcc/testsuite/c-c++-common/gomp/nesting-1.c
new file mode 100644
index 0000000..70b3e8d
--- /dev/null
+++ gcc/testsuite/c-c++-common/gomp/nesting-1.c
@@ -0,0 +1,77 @@
+void
+f_omp_parallel (void)
+{
+#pragma omp parallel
+  {
+    int i;
+
+#pragma omp parallel
+    ;
+
+#pragma omp target
+    ;
+
+#pragma omp target data
+    ;
+
+#pragma omp target update to(i)
+
+#pragma omp target data
+    {
+#pragma omp parallel
+      ;
+
+#pragma omp target
+      ;
+
+#pragma omp target data
+      ;
+
+#pragma omp target update to(i)
+    }
+  }
+}
+
+void
+f_omp_target (void)
+{
+#pragma omp target
+  {
+#pragma omp parallel
+    ;
+  }
+}
+
+void
+f_omp_target_data (void)
+{
+#pragma omp target data
+  {
+    int i;
+
+#pragma omp parallel
+    ;
+
+#pragma omp target
+    ;
+
+#pragma omp target data
+    ;
+
+#pragma omp target update to(i)
+
+#pragma omp target data
+    {
+#pragma omp parallel
+      ;
+
+#pragma omp target
+      ;
+
+#pragma omp target data
+      ;
+
+#pragma omp target update to(i)
+    }
+  }
+}
diff --git gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
new file mode 100644
index 0000000..6a65699
--- /dev/null
+++ gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
@@ -0,0 +1,23 @@
+void
+f_omp_target (void)
+{
+#pragma omp target
+  {
+    int i;
+
+#pragma omp target /* { dg-warning "target construct inside of target region" } */
+    ;
+#pragma omp target data /* { dg-warning "target data construct inside of target region" } */
+    ;
+#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */
+
+#pragma omp parallel
+    {
+#pragma omp target /* { dg-warning "target construct inside of target region" } */
+      ;
+#pragma omp target data /* { dg-warning "target data construct inside of target region" } */
+      ;
+#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */
+    }
+  }
+}


Grüße,
 Thomas

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