Re: [patch,gomp4] error on invalid acc loop clauses

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

On Fri, 15 May 2015 11:10:21 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch teaches the c and c++ front ends to error on invalid and
> conflicting acc loop clauses. E.g., an acc loop cannot have 'gang seq'
> and the worker and vector clauses inside parallel regions cannot have
> optional kernel-specific arguments.

Thanks!

> The c and c++ front end also error when it detects a parallel or kernels
> region nested inside a parallel or kernels region. E.g.
> 
>   #pragma acc parallel
>   {
>     #pragma acc parallel
>      ...
>   }

OK, but see below.

> I included two new test cases in this patch. They are mostly identical
> but, unfortunately, the c and c++ front ends emit slightly different
> error messages.

The preference is to keep these as single files (so that C and C++ can
easily be maintained together), and use the appropriate dg-* directives
to select the expected C or C++ error message, respectively, or use
regular expressions so as to match both the expected C and C++ error
variants in one go, if they're similar enough.

> The front ends still need to be cleaned before this functionality should
> be considered for mainline. So for the time being I've applied this
> patch to gomp-4_0-branch.

What remains to be done?

Then, what about the Fortran front end?  Checking already done as well as
test coverage existing, similar to C and C++?

Patch review comments:

> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -234,6 +234,10 @@ typedef struct GTY(()) c_parser {
>    /* True if we are in a context where the Objective-C "Property attribute"
>       keywords are valid.  */
>    BOOL_BITFIELD objc_property_attr_context : 1;
> +  /* True if we are inside a OpenACC parallel region.  */
> +  BOOL_BITFIELD oacc_parallel_region : 1;
> +  /* True if we are inside a OpenACC kernels region.  */
> +  BOOL_BITFIELD oacc_kernels_region : 1;

Hmm.

> @@ -10839,6 +10843,7 @@ c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind,
>  	  mark_exp_read (expr);
>  	  require_positive_expr (expr, expr_loc, str);
>  	  *op_to_parse = expr;
> +	  op_to_parse = &op0;
>  	}
>        while (!c_parser_next_token_is (parser, CPP_CLOSE_PAREN));
>        c_parser_consume_token (parser);
> @@ -10852,6 +10857,17 @@ c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind,
>    if (op1)
>      OMP_CLAUSE_OPERAND (c, 1) = op1;
>    OMP_CLAUSE_CHAIN (c) = list;
> +
> +  if (parser->oacc_parallel_region && (op0 != NULL || op1 != NULL))
> +    {
> +      if (c_kind != PRAGMA_OACC_CLAUSE_GANG)
> +	c_parser_error (parser, c_kind == PRAGMA_OACC_CLAUSE_WORKER ?
> +			"worker clause arguments are not supported in OpenACC parallel regions"
> +			: "vector clause arguments are not supported in OpenACC parallel regions");
> +      else if (op0 != NULL)
> +	c_parser_error (parser, "non-static argument to clause gang");
> +    }

Instead of in c_parser_oacc_shape_clause, shouldn't such checking rather
be done inside the function invoking c_parser_oacc_shape_clause, that is,
c_parser_oacc_parallel, etc.?

> @@ -12721,7 +12737,10 @@ static tree
>  c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
>  		    omp_clause_mask mask, tree *cclauses)
>  {
> -  tree stmt, clauses, block;
> +  tree stmt, clauses, block, c;
> +  bool gwv = false;
> +  bool auto_clause = false;
> +  bool seq_clause = false;
>  
>    strcat (p_name, " loop");
>    mask |= OACC_LOOP_CLAUSE_MASK;
> @@ -12732,6 +12751,33 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
>    if (cclauses)
>      clauses = oacc_split_loop_clauses (clauses, cclauses);
>  
> +  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
> +    {
> +      switch (OMP_CLAUSE_CODE (c))
> +	{
> +	case OMP_CLAUSE_GANG:
> +	case OMP_CLAUSE_WORKER:
> +	case OMP_CLAUSE_VECTOR:
> +	  gwv = true;
> +	  break;
> +	case OMP_CLAUSE_AUTO:
> +	  auto_clause = true;
> +	  break;
> +	case OMP_CLAUSE_SEQ:
> +	  seq_clause = true;
> +	  break;
> +	default:
> +	  ;
> +	}
> +    }
> +
> +  if (gwv && auto_clause)
> +    c_parser_error (parser, "incompatible use of clause %<auto%>");
> +  else if (gwv && seq_clause)
> +    c_parser_error (parser, "incompatible use of clause %<seq%>");
> +  else if (auto_clause && seq_clause)
> +    c_parser_error (parser, "incompatible use of clause %<seq%> and %<auto%>");
> +
>    block = c_begin_compound_stmt (true);
>    stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
>    block = c_end_compound_stmt (loc, block, true);

I would have expected such checking to be done in c_omp_finish_clauses --
But maybe it's also OK to do it here, given that the loop construct is
the only one where these clauses can appear.  Jakub, any strong
preference?

> @@ -12774,6 +12820,13 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
>  
>    strcat (p_name, " kernels");
>  
> +  if (parser->oacc_parallel_region || parser->oacc_kernels_region)
> +    {
> +      c_parser_error (parser, "nested kernels region");
> +    }
> +
> +  parser->oacc_kernels_region = true;

Regarding this...

> @@ -12843,6 +12898,13 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
>  
>    strcat (p_name, " parallel");
>  
> +  if (parser->oacc_parallel_region || parser->oacc_kernels_region)
> +    {
> +      c_parser_error (parser, "nested parallel region");
> +    }
> +
> +  parser->oacc_parallel_region = true;

..., and this: why not do such nesting checking in
gcc/omp-low.c:check_omp_nesting_restrictions?  Currently (changed by
Bernd in internal r442824, 2014-11-29) we're allowing all
OpenACC-inside-OpenACC nesting -- shouldn't that be changed instead of
repeating the checks in every front end (Jakub?)?

I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
Âgang, worker and vector may occur only once in a loop nestÂ, and Âgang,
worker and vector must occur in this order in a loop nestÂ.  Don't know
if that conceptually also belongs into
gcc/omp-low.c:check_omp_nesting_restrictions?

> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c

As for C.

> --- a/gcc/cp/parser.h
> +++ b/gcc/cp/parser.h

Likewise.

> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -2959,14 +2959,6 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
>  	    }
>  	  if (outer_type == GIMPLE_OMP_FOR)
>  	    outer_ctx->gwv_below |= val;
> -	  if (OMP_CLAUSE_OPERAND (c, 0) != NULL_TREE)
> -	    {
> -	      omp_context *enclosing = enclosing_target_ctx (outer_ctx);
> -	      if (gimple_omp_target_kind (enclosing->stmt)
> -		  == GF_OMP_TARGET_KIND_OACC_PARALLEL)
> -		error_at (gimple_location (stmt),
> -			  "no arguments allowed to gang, worker and vector clauses inside parallel");
> -	    }
>  	}
>      }

> --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
> @@ -7,9 +7,9 @@ f_acc_parallel (void)
>  {
>  #pragma acc parallel
>    {
> -#pragma acc parallel /* { dg-bogus "parallel construct inside of parallel region" "not implemented" { xfail *-*-* } } */
> +#pragma acc parallel /* { dg-error "nested parallel region" } */
>      ;
> -#pragma acc kernels /* { dg-bogus "kernels construct inside of parallel region" "not implemented" { xfail *-*-* } } */
> +#pragma acc kernels /* { dg-error " kernels construct inside of parallel" } */
>      ;
>  #pragma acc data /* { dg-error "data construct inside of parallel region" } */
>      ;
> @@ -26,9 +26,9 @@ f_acc_kernels (void)
>  {
>  #pragma acc kernels
>    {
> -#pragma acc parallel /* { dg-bogus "parallel construct inside of kernels region" "not implemented" { xfail *-*-* } } */
> +#pragma acc parallel /* { dg-error "nested parallel region" } */
>      ;
> -#pragma acc kernels /* { dg-bogus "kernels construct inside of kernels region" "not implemented" { xfail *-*-* } } */
> +#pragma acc kernels /* { dg-error "nested kernels region" } */
>      ;
>  #pragma acc data /* { dg-error "data construct inside of kernels region" } */
>      ;
> @@ -37,3 +37,8 @@ f_acc_kernels (void)
>  #pragma acc exit data delete(i) /* { dg-error "enter/exit data construct inside of kernels region" } */
>    }
>  }
> +
> +// { dg-error "parallel construct inside of parallel" "" { target *-*-* } 10 }
> +
> +// { dg-error "parallel construct inside of kernels" "" { target *-*-* } 29 }
> +// { dg-error "kernels construct inside of kernels" "" { target *-*-* } 31 }

I like it better if these are placed next to the line they apply to.
Makes it easier to see what's going on.

> --- /dev/null
> +++ b/gcc/testsuite/g++.dg/goacc/loop-4.C

You didn't actually commit this file.

> @@ -0,0 +1,686 @@
> +// { dg-do compile }
> +// { dg-options "-fopenacc" }
> +// { dg-additional-options "-fmax-errors=200" }

Inside gcc/testsuite/*/goacc/, -fopenacc is enabled by default.

> +
> +int
> +main ()
> +{
> +  int i, j;
> +
> +#pragma acc kernels
> +  {
> +[...]
> +#pragma acc loop gang(static:*)
> +    for (i = 0; i < 10; i++)
> +      { }
> +#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
> +    for (i = 0; i < 10; i++)
> +      {
> +[...]

Doesn't it make sense to separate test cases that are expected to be
handled successfully (no error messages) from those where we expect error
diagnostics, so that we can be sure the former ones will also be
successfully run through all later compiler passes?

> +[...]

> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 377 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 397 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 400 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 423 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 426 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 429 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 535 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 555 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 558 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 581 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 584 }
> +// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 587 }

Hmm, that doesn't happen for C?

> --- /dev/null
> +++ b/gcc/testsuite/gcc.dg/goacc/loop-1.c

You didn't actually commit this file.

> @@ -0,0 +1,674 @@
> +[...]
> +// { dg-excess-errors "invalid controlling predicate" }

Why is that needed?


GrÃÃe,
 Thomas

Attachment: pgpNQBnnmc0Qy.pgp
Description: PGP signature