On 05/20/2015 01:23 AM, Thomas Schwinge wrote:
>> 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?
Jakub made some general comments a couple of weeks ago when you applied
our internal changes to gomp-4_0-branch. I was planning on addressing
those comments first before requesting the merge to trunk.
> Then, what about the Fortran front end? Checking already done as well as
> test coverage existing, similar to C and C++?
Fortran is good.
> 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.
What's wrong with this? Fortran does something similar. Besides, this is
only temporary until OpenACC 2.5.
>> @@ -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.?
I don't think that will help. c_parser_oacc_shape_clause parses 'gang',
'worker' and 'vector' which aren't available to acc parallel or acc
kernels. Well, they are, be rigged up the front end to split acc
parallel loops and acc kernels loop.
>> @@ -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?)?
The fortran front end is doing this. Also, Joseph told me the front ends
should report error messages when possible. I have no problems reverting
back to the original behavior though.
> 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?
Yeah, someone needs to clean that up. I tried to keep this patch local
to the c and c++ front ends.
>> --- 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.
There's already a dg-error on those lines and I think dg-error can only
report one error per line.
>> --- /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.
True. However the toplevel dg.exp also executes those tests and it
doesn't use -fopenacc without dg-options.
>> +
>> +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?
I thought the ones that we expect to run successfully belong in libgomp?
>> +[...]
>
>> +// { 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?
I'm not sure why, unfortunately.
>> --- /dev/null
>> +++ b/gcc/testsuite/gcc.dg/goacc/loop-1.c
>
> You didn't actually commit this file.
Two files that I missed.
>> @@ -0,0 +1,674 @@
>> +[...]
>> +// { dg-excess-errors "invalid controlling predicate" }
>
> Why is that needed?
I'm not sure. I thought it had something to do with Bernd's predication
pass.
Cesar
