OpenACC for C++ front end

Hi!

On 11/13/2014 07:02 AM, Jakub Jelinek wrote:
> On Wed, Nov 05, 2014 at 03:37:08PM -0600, James Norris wrote:
>> 2014-11-05  James Norris  <jnorris@codesourcery.com>
>> 	    Cesar Philippidis  <cesar@codesourcery.com>
>> 	    Thomas Schwinge  <thomas@codesourcery.com>
>> 	    Ilmir Usmanov  <i.usmanov@samsung.com>
> ...
>
> Please check formatting.  I see various spots with 8 spaces instead of tabs,
> e.g. on
> +  check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH,
> +                                                "vector_length", location);
> even the alignment is wrong, I see calls without space before (:
> +  if (args == NULL || args->length() == 0)
> ...
> +             error("%<wait%> expression must be integral");
> other spots where the alignment isn't right:

Fixed.

> +static tree
> +cp_parser_oacc_cache (cp_parser *parser,
> +                               cp_token *pragma_tok __attribute__((unused)))
> (cp_token should be below cp_parser).  While at this,
> __attribute__((unused)) should be replaced by ATTRIBUTE_UNUSED, or removing
> the parameter name, or removing the parameter altogether even better.

Fixed by removing unused parameter.

> For the formatting issues, you can easily look for them in the patch
> (in lines starting with +), change the patch and apply interdiff to your
> tree.

Thank you for the pointer to interdiff!

>
>> --- a/gcc/c-family/c-omp.c
>> +++ b/gcc/c-family/c-omp.c
>> @@ -29,8 +29,47 @@ along with GCC; see the file COPYING3.  If not see
>>   #include "c-pragma.h"
>>   #include "gimple-expr.h"
>>   #include "langhooks.h"
>> +#include "omp-low.h"
> As Thomas? has said, you should include gomp-constants.h and use them in:
>
>> +    t = build_int_cst (integer_type_node, -2);  /* TODO: XXX FIX -2.  */
> spots like this.

Fixed.

>
>> --- a/gcc/c-family/c-pragma.c
>> +++ b/gcc/c-family/c-pragma.c
>> @@ -1180,6 +1180,16 @@ typedef struct
>>   static vec<pragma_ns_name> registered_pp_pragmas;
>>   
>>   struct omp_pragma_def { const char *name; unsigned int id; };
>> +static const struct omp_pragma_def oacc_pragmas[] = {
>> +  { "data", PRAGMA_OACC_DATA },
>> +  { "enter", PRAGMA_OACC_ENTER_DATA },
>> +  { "exit", PRAGMA_OACC_EXIT_DATA },
>> +  { "kernels", PRAGMA_OACC_KERNELS },
>> +  { "loop", PRAGMA_OACC_LOOP },
>> +  { "parallel", PRAGMA_OACC_PARALLEL },
>> +  { "update", PRAGMA_OACC_UPDATE },
>> +  { "wait", PRAGMA_OACC_WAIT },
> I'd avoid the , after the last element.

Fixed.

>
>> @@ -1383,6 +1402,17 @@ c_invoke_pragma_handler (unsigned int id)
>>   void
>>   init_pragma (void)
>>   {
>> +  if (flag_openacc)
>> +    {
>> +      const int n_oacc_pragmas
>> +	= sizeof (oacc_pragmas) / sizeof (*oacc_pragmas);
>> +      int i;
>> +
>> +      for (i = 0; i < n_oacc_pragmas; ++i)
>> +	cpp_register_deferred_pragma (parse_in, "acc", oacc_pragmas[i].name,
>> +				      oacc_pragmas[i].id, true, true);
>> +    }
>> +
>>     if (flag_openmp)
>>       {
>>         const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
> Is -fopenmp -fopenacc tested not to run out of number of supported pragmas
> by libcpp?

Tests will be added in a follow-up patch once the middle end has been 
accepted.

>
>> @@ -65,23 +74,30 @@ typedef enum pragma_kind {
>>   } pragma_kind;
>>   
>>   
>> -/* All clauses defined by OpenMP 2.5, 3.0, 3.1 and 4.0.
>> +/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, and 4.0.
>>      Used internally by both C and C++ parsers.  */
>>   typedef enum pragma_omp_clause {
>>     PRAGMA_OMP_CLAUSE_NONE = 0,
>>   
>>     PRAGMA_OMP_CLAUSE_ALIGNED,
>> +  PRAGMA_OMP_CLAUSE_ASYNC,
>>     PRAGMA_OMP_CLAUSE_COLLAPSE,
>> +  PRAGMA_OMP_CLAUSE_COPY,
>>     PRAGMA_OMP_CLAUSE_COPYIN,
>> +  PRAGMA_OMP_CLAUSE_COPYOUT,
>>     PRAGMA_OMP_CLAUSE_COPYPRIVATE,
>> +  PRAGMA_OMP_CLAUSE_CREATE,
>>     PRAGMA_OMP_CLAUSE_DEFAULT,
>> +  PRAGMA_OMP_CLAUSE_DELETE,
>>     PRAGMA_OMP_CLAUSE_DEPEND,
>>     PRAGMA_OMP_CLAUSE_DEVICE,
>> +  PRAGMA_OMP_CLAUSE_DEVICEPTR,
>>     PRAGMA_OMP_CLAUSE_DIST_SCHEDULE,
>>     PRAGMA_OMP_CLAUSE_FINAL,
>>     PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
>>     PRAGMA_OMP_CLAUSE_FOR,
>>     PRAGMA_OMP_CLAUSE_FROM,
>> +  PRAGMA_OMP_CLAUSE_HOST,
>>     PRAGMA_OMP_CLAUSE_IF,
>>     PRAGMA_OMP_CLAUSE_INBRANCH,
>>     PRAGMA_OMP_CLAUSE_LASTPRIVATE,
>> @@ -90,16 +106,24 @@ typedef enum pragma_omp_clause {
>>     PRAGMA_OMP_CLAUSE_MERGEABLE,
>>     PRAGMA_OMP_CLAUSE_NOTINBRANCH,
>>     PRAGMA_OMP_CLAUSE_NOWAIT,
>> +  PRAGMA_OMP_CLAUSE_NUM_GANGS,
>>     PRAGMA_OMP_CLAUSE_NUM_TEAMS,
>>     PRAGMA_OMP_CLAUSE_NUM_THREADS,
>> +  PRAGMA_OMP_CLAUSE_NUM_WORKERS,
>>     PRAGMA_OMP_CLAUSE_ORDERED,
>>     PRAGMA_OMP_CLAUSE_PARALLEL,
>> +  PRAGMA_OMP_CLAUSE_PRESENT,
>> +  PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
>> +  PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
>> +  PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT,
>> +  PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE,
>>     PRAGMA_OMP_CLAUSE_PRIVATE,
>>     PRAGMA_OMP_CLAUSE_PROC_BIND,
>>     PRAGMA_OMP_CLAUSE_REDUCTION,
>>     PRAGMA_OMP_CLAUSE_SAFELEN,
>>     PRAGMA_OMP_CLAUSE_SCHEDULE,
>>     PRAGMA_OMP_CLAUSE_SECTIONS,
>> +  PRAGMA_OMP_CLAUSE_SELF,
>>     PRAGMA_OMP_CLAUSE_SHARED,
>>     PRAGMA_OMP_CLAUSE_SIMDLEN,
>>     PRAGMA_OMP_CLAUSE_TASKGROUP,
>> @@ -107,6 +131,8 @@ typedef enum pragma_omp_clause {
>>     PRAGMA_OMP_CLAUSE_TO,
>>     PRAGMA_OMP_CLAUSE_UNIFORM,
>>     PRAGMA_OMP_CLAUSE_UNTIED,
>> +  PRAGMA_OMP_CLAUSE_VECTOR_LENGTH,
>> +  PRAGMA_OMP_CLAUSE_WAIT,
> Like for CILK, I'd strongly prefer if for the clauses that are
> specific to OpenACC only you'd use PRAGMA_OACC_CLAUSE_* instead,
> and put them after the PRAGMA_CILK_* enum values.
> If you want to have PRAGMA_OACC_CLAUSE_ aliases also for the
> clauses shared in between OpenMP and OpenACC, feel free to add
> aliases like Cilk+ has them.

Fixed by renaming OpenACC-specific clauses. Have deferred
on creating alias names.

>    It is unfortunately lots of new clauses
> and we are getting close to the 64 clauses limit :( when they are used
> in bitmasks.
>

Will put my 'thinking cap' on.

>> --- a/gcc/cp/parser.c
>> +++ b/gcc/cp/parser.c
>> @@ -27437,6 +27437,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
>>       result = PRAGMA_OMP_CLAUSE_IF;
>>     else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DEFAULT))
>>       result = PRAGMA_OMP_CLAUSE_DEFAULT;
>> +  else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DELETE))
>> +    result = PRAGMA_OMP_CLAUSE_DELETE;
>>     else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_PRIVATE))
>>       result = PRAGMA_OMP_CLAUSE_PRIVATE;
>>     else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
>> @@ -27451,20 +27453,30 @@ cp_parser_omp_clause_name (cp_parser *parser)
>>   	case 'a':
>>   	  if (!strcmp ("aligned", p))
>>   	    result = PRAGMA_OMP_CLAUSE_ALIGNED;
>> +	  else if (!strcmp ("async", p))
>> +	    result = PRAGMA_OMP_CLAUSE_ASYNC;
> Please adjust to PRAGMA_OACC_CLAUSE_* here etc. too; that will make it
> more clear that the clauses are OpenACC specific.

Fix as described above.

>> +      /* FIXME diagnostics: Ideally we should keep individual
>> +	 locations for all the variables in the var list to make the
>> +	 following errors more precise.  Perhaps
>> +	 c_parser_omp_var_list_parens should construct a list of
>> +	 locations to go along with the var list.  */
> Please avoid writing todo lists into the source, that should be
> tracked in some PR instead.  It is a general thing, we really should
> have some expression that would be used in the FEs
> to have a FE only expression code that would wrap decls and constants
> that don't have a location to provide the location for them, then
> we can use it everywhere, including the omp variable lists.
> We need it for better diagnostics (to have right locus or locus ranges
> in FE diagnostics).

Will submit a PR.

>> +  check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH,
>> +                                                "vector_length", location);
> See above about formatting.

Fixed.

>
>> +static tree
>> +cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list)
>> +{
>> +  vec<tree, va_gc> *args;
>> +  tree t, args_tree;
>> +
>> +  args = cp_parser_parenthesized_expression_list (parser, non_attr,
>> +                                                  /*cast_p=*/false,
>> +                                                  /*allow_expansion_p=*/true,
>> +                                                  /*non_constant_p=*/NULL);
>> +
>> +  if (args == NULL || args->length() == 0)
>> +    {
>> +      cp_parser_error (parser, "expected integer expression before ')'");
> Is cp_parser_parenthesized_expression_list an integer expression?  Sounds
> the wording doesn't match the implementation.  Also, please use
> %<)%> instead of ')'.

Fixed.

>> +      tree targ = TREE_VALUE (t);
>> +
>> +      if (targ != error_mark_node)
>> +        {
>> +	  if (!INTEGRAL_TYPE_P (TREE_TYPE (targ)))
>> +	    {
>> +	      error("%<wait%> expression must be integral");
>> +	      targ = error_mark_node;
> Why do you assign to targ at all, when it is not used afterwards?
> If you remove it, drop the {}s around error too.  And see above for
> formatting comment.

Fixed.

>> +  if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
>> +      || !tree_fits_shwi_p (t)
>> +      || (n = tree_to_shwi (t)) <= 0
>> +      || (int) n != n)
>> +    {
>> +      error_at (location, "expected positive integer expression");
>> +      return list;
>> +    }
> Do you use it as int value anywhere, or only pass through as tree to some
> function?  If the latter, then neither the tree_fits* nor (int) n != n check
> is needed, all you should care about is INTEGRAL_TYPE_P and tree_int_cst_sgn
>> 0.  Also, does it really have to be INTEGER_CST already during parsing, or
> is it enough if it is INTEGER_CST after folding/template instantiation etc.?
>
> The reason why omp_clause_collapse does the above is that the integer is
> directly used during the parsing.  For other integers, it is usually checked
> in finish_omp_clauses (after instantiation) or so.
>
>> +  if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
>> +      || !tree_fits_shwi_p (t)
>> +      || (n = tree_to_shwi (t)) <= 0
>> +      || (int) n != n)
>> +    {
>> +      error_at (location, "expected positive integer expression");
>> +      return list;
>> +    }
> Likewise.

Fixed both occurences.

>
>> +static tree
>> +cp_parser_oacc_clause_async (cp_parser *parser, tree list)
>> +{
>> +  tree c, t;
>> +  location_t loc = cp_lexer_peek_token (parser->lexer)->location;
>> +
>> +  /* TODO XXX: FIX -1  (acc_async_noval).  */
> Again, please use gomp-constants.h constant here.
>
>> +  t = build_int_cst (integer_type_node, -1);

Fixed.

>> +/* OpenACC 2.0:
>> +   # pragma acc enter data oacc-enter-data-clause[optseq] new-line
>> +
>> +   or
>> +
>> +   # pragma acc exit data oacc-exit-data-clause[optseq] new-line
>> +
>> +   LOC is the location of the #pragma token.
> Don't understand the comment here, LOC isn't mentioned anywhere in there,
> nor it is a parameter of the function.

Comment removed along with similar occurences.

>
>> +
>> +  if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL)
>> +     || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME))
>> +    {
>> +      cp_parser_error (parser, enter
>> +		       ? "expected %<data%> in %<#pragma acc enter data%>"
>> +		       : "expected %<data%> in %<#pragma acc exit data%>");
> Shouldn't the diagnostics for
> #pragma acc enter
> #pragma acc enter for
> #pragma acc enter data2
> be all the same?  The cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL)
> test doesn't make any sense together with
> || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME), for CPP_PRAGMA_EOL
> the next token is not CPP_NAME.
> Also, as #pragma acc enter data is the name of the directive, it shouldn't
> be translated, so IMHO you can just use
> "expected %<data%> after %<#pragma acc %s%>", enter ? "enter" : "exit".

Fixed.

>
>> +      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
>> +      return NULL_TREE;
>> +    }
>> +
>> +  const char *p =
>> +    IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
>> +  if (strcmp (p, "data") != 0)
>> +    {
>> +      cp_parser_error (parser, "invalid pragma");
> As written above, best set p to "", if
> cp_lexer_next_token_is (parser->lexer, CPP_NAME) is true, set it to
> IDENTIFIER_POINTER ... and then strcmp  and use one diagnostics for that.

Fixed.

>> +  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
>> +    {
>> +      error_at (pragma_tok->location,
>> +		"%<#pragma acc enter data%> has no data movement clause");
> Unconditional enter even when it could be exit?

Fixed.

>
>> +	case OMP_CLAUSE_ASYNC:
> Wonder if similar thing to PRAGMA_OACC_CLAUSE_* shouldn't be used here,
> either OACC_CLAUSE_ASYNC or OMP_CLAUSE_ACC_ASYNC to make it clear
> it is OpenACC thing.  Though in this case I don't feel as strong about it as
> in case of PRAGMA_OACC_CLAUSE_*.

I don't have a strong feeling one way or another either, so I'll
leave the OMP_CLAUSES_* alone for the time being.

>
>
>> +	  t = OMP_CLAUSE_ASYNC_EXPR (c);
>> +	  if (t == error_mark_node)
>> +	    remove = true;
>> +	  else if (!type_dependent_expression_p (t)
>> +		   && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
>> +	    {
>> +	      error ("%<async%> expression must be integral");
> You have OMP_CLAUSE_LOCATION (c) which you could use for error_at.

I followed the convention that was used elsewhere in the function
at using error ().

>
> Otherwise LGTM.

Thank you for taking the time to review!

OK to commit after middle end has been accepted?

Regards,
Jim
=> cp/ChangeLog

2014-XX-XX  James Norris  <jnorris@codesourcery.com>
        Cesar Philippidis  <cesar@codesourcery.com>
        Thomas Schwinge  <thomas@codesourcery.com>
        Ilmir Usmanov  <i.usmanov@samsung.com>

    * cp-tree.h (finish_oacc_data, finish_oacc_kernels,
    finish_oacc_parallel): New prototypes.
    * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
    OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT.
    (finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
    functions.
    * parser.c (cp_parser_omp_clause_name): Don't parse the
    identifier for RID_DELETE. Add parsing of OpenACC clauses.
    (cp_parser_omp_var_list_no_open): Add handling of array specifier.
    (cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr,
    cp_parser_oacc_vector_length, cp_parser_oacc_wait_list,
    cp_parser_oacc_clause_wait, cp_parser_omp_clause_num_gangs,
    cp_parser_omp_clause_num_workers, cp_parser_oacc_clause_async,
    cp_parser_oacc_all_clauses, cp_parser_oacc_cache, cp_parser_oacc_data,
    cp_parser_oacc_enter_exit_data, cp_parser_oacc_kernels,
    cp_parser_oacc_loop, cp_parser_oacc_parallel, cp_parser_oacc_update,
    cp_parser_oacc_wait): New functions.
    (cp_parser_omp_construct): Handle OpenACC pragmas.
    (cp_parser_pragma): Likewise.

    => c-family/ChangeLog

2014-XX-XX  James Norris  <jnorris@codesourcery.com>
        Cesar Philippidis  <cesar@codesourcery.com>
        Thomas Schwinge  <thomas@codesourcery.com>
        Ilmir Usmanov  <i.usmanov@samsung.com>

    * c-pragma.h (pragma_kind): Add OpenACC pragmas.
    (pragma_omp_clause): Add OpenACC clauses.
    * c-pragma.c (oacc_pragmas): New pragma definition array.
    (c_pp_lookup_pragma, init_pragma): Handle OpenACC pragmas.
    * c-cppbuiltin.c (c_cpp_builtins): Conditionally define _OPENACC.
    * c.opt (fopenacc): New option.
    * c-omp.c (c_finish_oacc_wait): New function.
    (c_omp_split_clauses): Catch OACC_PARALLEL.
    * c-common.h (c_finish_oacc_wait): New prototype.
    * c-common.c (DEF_FUNCTION_TYPE_8, DEF_FUNTION_TYPE_12): New defines.

OpenACC for C++ front end

Commit Message

Comments

Patch