[PATCH] OpenACC for C front end

[James Norris]

Hi!

On 11/13/2014 09:04 AM, Jakub Jelinek wrote:
> On Wed, Nov 05, 2014 at 03:39:44PM -0600, James Norris wrote:
>> 	* c-typeck.c (c_finish_oacc_parallel, c_finish_oacc_kernels,
>> 	c_finish_oacc_data): New functions.
>> 	(handle_omp_array_sections, c_finish_omp_clauses):
> Handle should be on the above line, no need to wrap too early.

Fixed.

>
>> @@ -9763,6 +9830,10 @@ c_parser_omp_clause_name (c_parser *parser)
>>   	  else if (!strcmp ("from", p))
>>   	    result = PRAGMA_OMP_CLAUSE_FROM;
>>   	  break;
>> +	case 'h':
>> +	  if (!strcmp ("host", p))
>> +	    result = PRAGMA_OMP_CLAUSE_SELF;
>> +	  break;
> Shouldn't this be PRAGMA_OMP_CLAUSE_HOST (PRAGMA_OACC_CLAUSE_HOST)
> instead?  It is _HOST in the C++ patch, are there no C tests with
> that clause covering it?

The "host" clause is a synonym for the "self" clause. The initial
C++ patch did not treat "host" as a synonym and has amended
accordingly.

>> +      tree v = TREE_PURPOSE (t);
>> +
>> +      /* 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.  */
> Like in C++ patch, please avoid the comment, file a PR instead,
> or just queue the work for GCC 6.

Will submit a PR.

>> +      /* Attempt to statically determine when the number isn't positive.  */
>> +      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
>> +		       build_int_cst (TREE_TYPE (t), 0));
> build_int_cst not aligned below expr_loc.

Fixed.

>> +      if (CAN_HAVE_LOCATION_P (c))
>> +	SET_EXPR_LOCATION (c, expr_loc);
>> +      if (c == boolean_true_node)
>> +	{
>> +	  warning_at (expr_loc, 0,
>> +		      "%<num_gangs%> value must be positive");
> This would fit perfectly on one line.

Fixed.

>> +  tree c, t;
>> +  location_t loc = c_parser_peek_token (parser)->location;
>> +
>> +  /* TODO XXX: FIX -1  (acc_async_noval).  */
>> +  t = build_int_cst (integer_type_node, -1);
> Again, as in C++ patch, please avoid this.  Use gomp-constants.h,
> or some enum, or explain what the -1 is, but avoid TODO XXX: FIX.

Fixed.

>> +  else
>> +    {
>> +      t = c_fully_fold (t, false, NULL);
>> +    }
> Please avoid the {}s and reindent.

Fixed.

>
>> -/* OpenMP 4.0:
>> -   parallel
>> -   for
>> -   sections
>> -   taskgroup */
>> +      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
>> +	{
>> +	  c_parser_error (parser, "expected integer expression");
>> +	  return list;
>> +	}
>>   
>> -static tree
>> -c_parser_omp_clause_cancelkind (c_parser *parser ATTRIBUTE_UNUSED,
>> -				enum omp_clause_code code, tree list)
>> -{
>> -  tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
>> +      /* Attempt to statically determine when the number isn't positive.  */
>> +      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
>> +		       build_int_cst (TREE_TYPE (t), 0));
> I wonder if it wouldn't be better to just put the new OpenACC routines
> into a new block of code, not stick it in between the OpenMP handling
> routines, because diff apparently lost track and I'm afraid so will svn blame/git blame
> and we'll lose history for the OpenMP changes.

There was a mistake in naming the function: 
c_parser_omp_clause_vector_length.
Once it was renamed to: c_parser_oacc_clause_vector_length, diff was able to
keep track.

>> +	case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH:
>> +	  clauses = c_parser_omp_clause_vector_length (parser, clauses);
>> +	  c_name = "vector_length";
>> +	  break;
> That is OpenACC clause, right?  Shouldn't the routine be called
> c_parser_oacc_clause_vector_length?

Fixed.

>> +	case PRAGMA_OMP_CLAUSE_WAIT:
>> +	  clauses = c_parser_oacc_clause_wait (parser, clauses);
> E.g. c_parser_oacc_clause_wait is.

Fixed.

>
>> +  clauses =  c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
>> +					"#pragma acc data");
> Too many spaces after =.

Fixed.

>> +/* OpenACC 2.0:
>> +   # pragma acc kernels oacc-kernels-clause[optseq] new-line
>> +     structured-block
>> +
>> +   LOC is the location of the #pragma token.
> Again, what is LOC?

Fixed by removal of comment, along with other occurences.

>> +  clauses =  c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
>> +					p_name);
> See above.

Fixed.

>> +      c_parser_error (parser, enter
>> +		      ? "expected %<data%> in %<#pragma acc enter data%>"
>> +		      : "expected %<data%> in %<#pragma acc exit data%>");
>> +      c_parser_skip_to_pragma_eol (parser);
>> +      return;
>> +    }
>> +
>> +  const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
>> +  if (strcmp (p, "data") != 0)
>> +    {
>> +      c_parser_error (parser, "invalid pragma");
> See the C++ patch review.

Fixed.

>> +  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
>> +    {
>> +      error_at (loc, enter
>> +		? "%<#pragma acc enter data%> has no data movement clause"
>> +		: "%<#pragma acc exit data%> has no data movement clause");
> Similarly (though, in this case C++ had unconditional acc enter data).

Fixed.

>
>> +  clauses =  c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
>> +					p_name);
> See above.

Fixed.

>
>> +  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
>> +    {
>> +      error_at (loc,
>> +		"%<#pragma acc update%> must contain at least one "
>> +		"%<device%> or %<host/self%> clause");
> There is no host/self clause, so you better write or %<host%> or
> %<self%> clause?

Fixed.

>
> Otherwise LGTM.

Thank you for taking the time to review!

OK to commit after middle end is accepted?

Regards,
Jim

-------------- next part --------------
    => c/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-typeck.c (c_finish_oacc_parallel, c_finish_oacc_kernels,
    c_finish_oacc_data): New functions.
    (handle_omp_array_sections): Catch MAP_FORCE_DEVICEPTR.
    (c_finish_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR,
    OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH,
    OMP_CLAUSE_ASYNC, and OMP_CLAUSE_WAIT.
    * c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels,
    c_finish_oacc_data): New prototypes.
    * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA,
    PRAGMA_OACC_EXIT_DATA, and PRAGMA_OACC_UPDATE.
    (c_parser_omp_clause_name): Handle OpenACC clauses.
    (c_parser_oacc_wait_list, c_parser_oacc_data_clause,
    c_parser_oacc_data_clause_deviceptr, c_parser_omp_clause_num_gangs,
    c_parser_omp_clause_num_workers, c_parser_oacc_clause_async,
    c_parser_oacc_clause_wait, c_parser_omp_clause_vector_length,
    c_parser_oacc_all_clauses, c_parser_oacc_cache, c_parser_oacc_data,
    c_parser_oacc_kernels, c_parser_oacc_enter_exit_data, c_parser_oacc_loop,
    c_parser_oacc_parallel, c_parser_oacc_update, c_parser_oacc_wait):
    New functions.
    (c_parser_omp_construct): Handle PRAGMA_OACC_CACHE, PRAGMA_OMP_DATA,
    PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL and
    PRAMGA_OACC_WAIT.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: new2.diff
Type: text/x-patch
Size: 38880 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20141121/6bf74827/attachment.bin>