Re: [PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end

[Thomas Schwinge <thomas@codesourcery.com>]

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

Hi!

On Fri, 31 Jan 2014 15:16:07 +0400, Ilmir Usmanov <i.usmanov@samsung.com> wrote:
>      OpenACC 1.0 support -- GENERIC nodes and gimplify stubs.

Thanks!  This one is nearly ready for commit, and can then go in already,
while the Fortran front end patches are still being dicussed.  :-)

Please merge »OpenACC 1.0 support -- documentation« into this patch.  In
gcc/doc/generic.texi, please also add an entry for OACC_DECLARE, which is
currently missing.  Ideally, gcc/doc/generic.texi <OMP_CLAUSE> should
also be updated for the OpenACC clauses (as well as recently added OpenMP
ones), but as that one's outdated already, this is not a prerequisite.

For ChangeLog files updates (on gomp-4_0-branch, use the respective
ChangeLog.gomp files, by the way), should just you be listed as the
author, or also your colleagues?

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c

> @@ -6157,6 +6166,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	    remove = true;
>  	  break;
>  
> +  case OMP_CLAUSE_HOST:
> +  case [...]
> +  case OMP_CLAUSE_VECTOR_LENGTH:
>  	case OMP_CLAUSE_NOWAIT:
>  	case OMP_CLAUSE_ORDERED:
>  	case OMP_CLAUSE_UNTIED:

Indentation (one tab instead of two spaces).

> @@ -6476,6 +6499,23 @@ gimplify_adjust_omp_clauses (tree *list_p)
>  	    }
>  	  break;
>  
> +  case OMP_CLAUSE_HOST:
> +  case [...]
> +  case OMP_CLAUSE_VECTOR_LENGTH:
> +    sorry ("Clause not supported yet");
> +    break;
> +
>  	case OMP_CLAUSE_REDUCTION:
>  	case OMP_CLAUSE_COPYIN:
>  	case OMP_CLAUSE_COPYPRIVATE:

Indentation.

Also, shouldn't the sorry be moved to gimplify_scan_omp_clauses, and a
»remove = true« be added in there, and gimplify_adjust_omp_clauses then
just contain a gcc_unreachable (or, move the new »case OMP_CLAUSE_*« next
to the existing default: that already contains gcc_unreachable)?

> @@ -7988,6 +8028,19 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
>  	  ret = GS_ALL_DONE;
>  	  break;
>  
> +  case OACC_KERNELS:
> +  case OACC_DATA:
> +  case OACC_CACHE:
> +  case OACC_WAIT:
> +  case OACC_HOST_DATA:
> +  case OACC_DECLARE:
> +  case OACC_UPDATE:
> +  case OACC_ENTER_DATA:
> +  case OACC_EXIT_DATA:
> +    sorry ("directive not yet implemented");
> +    ret = GS_ALL_DONE;
> +    break;
> +
>  	case OMP_PARALLEL:
>  	  gimplify_omp_parallel (expr_p, pre_p);
>  	  ret = GS_ALL_DONE;

Indentation.

Further down in gimplify_expr, shouldn't these new OACC_* codes also be
added to the »These expressions should already be in gimple IR form«
assert?

> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -1491,6 +1491,18 @@ fixup_child_record_type (omp_context *ctx)
>    TREE_TYPE (ctx->receiver_decl) = build_pointer_type (type);
>  }
>  
> +static bool
> +gimple_code_is_oacc (const_gimple g)
> +{
> +  switch (gimple_code (g))
> +    {
> +    case GIMPLE_OACC_PARALLEL:
> +      return true;
> +    default:
> +      return false;
> +    }
> +}
> +

Eventually, this will probably end up next to CASE_GIMPLE_OP/is_gimple_op
in gimple.h (or the latter be reworked to be able to ask for is_omp vs.
is_oacc vs. is_omp_or_oacc), but it's fine to do that once we actually
need it in files other than just omp-low.c, and once we support more
GIMPLE_OACC_* codes.

> @@ -1710,17 +1732,34 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)

> +  case OMP_CLAUSE_HOST:
> +  case OMP_CLAUSE_OACC_DEVICE:
> +  case OMP_CLAUSE_DEVICE_RESIDENT:
> +  case OMP_CLAUSE_USE_DEVICE:
> +  case OMP_CLAUSE_ASYNC:
> +  case OMP_CLAUSE_GANG:
> +  case OMP_CLAUSE_WAIT:
> +  case OMP_NO_CLAUSE_CACHE:
> +  case OMP_CLAUSE_INDEPENDENT:
> +  case OMP_CLAUSE_WORKER:
> +  case OMP_CLAUSE_VECTOR:
> +  case OMP_CLAUSE_NUM_GANGS:
> +  case OMP_CLAUSE_NUM_WORKERS:
> +  case OMP_CLAUSE_VECTOR_LENGTH:
> +    sorry ("Clause not supported yet");
> +    break;
> +
>  	default:
>  	  gcc_unreachable ();
>  	}

Indentation.

> @@ -1827,9 +1876,26 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
>  	case OMP_CLAUSE__LOOPTEMP_:
>  	case OMP_CLAUSE_TO:
>  	case OMP_CLAUSE_FROM:
> -	  gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
> +	  gcc_assert (!gimple_code_is_oacc (ctx->stmt));
>  	  break;
>  
> +  case OMP_CLAUSE_HOST:
> +  case OMP_CLAUSE_OACC_DEVICE:
> +  case OMP_CLAUSE_DEVICE_RESIDENT:
> +  case OMP_CLAUSE_USE_DEVICE:
> +  case OMP_CLAUSE_ASYNC:
> +  case OMP_CLAUSE_GANG:
> +  case OMP_CLAUSE_WAIT:
> +  case OMP_NO_CLAUSE_CACHE:
> +  case OMP_CLAUSE_INDEPENDENT:
> +  case OMP_CLAUSE_WORKER:
> +  case OMP_CLAUSE_VECTOR:
> +  case OMP_CLAUSE_NUM_GANGS:
> +  case OMP_CLAUSE_NUM_WORKERS:
> +  case OMP_CLAUSE_VECTOR_LENGTH:
> +    sorry ("Clause not supported yet");
> +    break;
> +
>  	default:
>  	  gcc_unreachable ();
>  	}

Indentation.

> --- a/gcc/tree-core.h
> +++ b/gcc/tree-core.h
> @@ -213,19 +213,19 @@ enum omp_clause_code {
>       (c_parser_omp_variable_list).  */
>    OMP_CLAUSE_ERROR = 0,
>  
> -  /* OpenMP clause: private (variable_list).  */
> +  /* OpenMP/OpenACC clause: private (variable_list).  */
>    OMP_CLAUSE_PRIVATE,
>  
>    /* OpenMP clause: shared (variable_list).  */
>    OMP_CLAUSE_SHARED,
>  
> -  /* OpenMP clause: firstprivate (variable_list).  */
> +  /* OpenMP/OpenACC clause: firstprivate (variable_list).  */
>    OMP_CLAUSE_FIRSTPRIVATE,
>  
>    /* OpenMP clause: lastprivate (variable_list).  */
>    OMP_CLAUSE_LASTPRIVATE,
>  
> -  /* OpenMP clause: reduction (operator:variable_list).
> +  /* OpenMP/OpenACC clause: reduction (operator:variable_list).
>       OMP_CLAUSE_REDUCTION_CODE: The tree_code of the operator.
>       Operand 1: OMP_CLAUSE_REDUCTION_INIT: Stmt-list to initialize the var.
>       Operand 2: OMP_CLAUSE_REDUCTION_MERGE: Stmt-list to merge private var
> @@ -265,10 +265,40 @@ enum omp_clause_code {
>       OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list).  */
>    OMP_CLAUSE_MAP,
>  
> +  /* OpenACC clause: host (variable_list).  */
> +  OMP_CLAUSE_HOST,
> +
> +  /* OpenACC clause: device (variable_list).  */
> +  OMP_CLAUSE_OACC_DEVICE,
> +
> +  /* OpenACC clause: device_resident (variable_list).  */
> +  OMP_CLAUSE_DEVICE_RESIDENT,
> +
> +  /* OpenACC clause: use_device (variable_list).  */
> +  OMP_CLAUSE_USE_DEVICE,
> +
> +  /* OpenACC clause: async [(integer-expression)].  */
> +  OMP_CLAUSE_ASYNC,
> +
> +  /* OpenACC clause: gang [(gang-argument-list)]. 
> +     Where 
> +      gang-argument-list: [gang-argument-list, ] gang-argument 
> +      gang-argument: [num:] integer-expression
> +                   | static: size-expression
> +      size-expression: * | integer-expression.  */
> +  OMP_CLAUSE_GANG,
> +
> +  /* OpenACC clause/directive: wait [(integer-expression-list)].  */
> +  OMP_CLAUSE_WAIT,
> +
> +  /* Internal structure to hold OpenACC cache directive's variable-list.
> +     #pragma acc cache (variable-list).  */
> +  OMP_NO_CLAUSE_CACHE,
> +
>    /* Internal clause: temporary for combined loops expansion.  */
>    OMP_CLAUSE__LOOPTEMP_,
>  
> -  /* OpenMP clause: if (scalar-expression).  */
> +  /* OpenMP/OpenACC clause: if (scalar-expression).  */
>    OMP_CLAUSE_IF,
>  
>    /* OpenMP clause: num_threads (integer-expression).  */
> @@ -281,12 +311,13 @@ enum omp_clause_code {
>    OMP_CLAUSE_NOWAIT,
>  
>    /* OpenMP clause: ordered.  */
> +  /* OpenACC clause: seq.  */
>    OMP_CLAUSE_ORDERED,
>  
>    /* OpenMP clause: default.  */
>    OMP_CLAUSE_DEFAULT,
>  
> -  /* OpenMP clause: collapse (constant-integer-expression).  */
> +  /* OpenMP/OpenACC clause: collapse (constant-integer-expression).  */
>    OMP_CLAUSE_COLLAPSE,
>  
>    /* OpenMP clause: untied.  */
> @@ -338,7 +369,25 @@ enum omp_clause_code {
>    OMP_CLAUSE_TASKGROUP,
>  
>    /* Internally used only clause, holding SIMD uid.  */
> -  OMP_CLAUSE__SIMDUID_
> +  OMP_CLAUSE__SIMDUID_,
> +
> +  /* OpenACC clause: independent.  */
> +  OMP_CLAUSE_INDEPENDENT,
> +
> +  /* OpenACC clause: worker [( [num:] integer-expression)].  */
> +  OMP_CLAUSE_WORKER,
> +
> +  /* OpenACC clause: vector [( [length:] integer-expression)].  */
> +  OMP_CLAUSE_VECTOR,
> +
> +  /* OpenACC clause: num_gangs (integer-expression).  */
> +  OMP_CLAUSE_NUM_GANGS,
> +
> +  /* OpenACC clause: num_workers (integer-expression).  */
> +  OMP_CLAUSE_NUM_WORKERS,
> +
> +  /* OpenACC clause: vector_length (integer-expression).  */
> +  OMP_CLAUSE_VECTOR_LENGTH
>  };

Unless I'm confused, some of the new clauses (those that deal with more
than just a single operand) will need to get special handling added à la
the other omp_clause_subcode variants in tree-core.h:tree_omp_clause?
(It is fine to rework that in a later patch, once we actually implement
these clauses.)

> --- a/gcc/tree-pretty-print.c
> +++ b/gcc/tree-pretty-print.c
> @@ -326,6 +326,21 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)

OMP_CLAUSE_WAIT is missing.

> @@ -2384,6 +2452,41 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,

OACC_ENTER_DATA and OACC_EXIT_DATA are missing.

> --- a/gcc/tree.c
> +++ b/gcc/tree.c

> @@ -327,7 +350,13 @@ const char * const omp_clause_code_name[] =

> +  "indepentend",

Typo: independent.

> @@ -11034,6 +11063,18 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,

OMP_CLAUSE_GANG missing.

>      case OMP_CLAUSE:
>        switch (OMP_CLAUSE_CODE (*tp))
>  	{
> +  case OMP_CLAUSE_HOST:
> +  case OMP_CLAUSE_OACC_DEVICE:
> +  case OMP_CLAUSE_DEVICE_RESIDENT:
> +  case OMP_CLAUSE_USE_DEVICE:
> +  case OMP_NO_CLAUSE_CACHE:
> +  case OMP_CLAUSE_ASYNC:
> +  case OMP_CLAUSE_WORKER:
> +  case OMP_CLAUSE_VECTOR:
> +  case OMP_CLAUSE_NUM_GANGS:
> +  case OMP_CLAUSE_NUM_WORKERS:
> +  case OMP_CLAUSE_VECTOR_LENGTH:
> +  case OMP_CLAUSE_WAIT:
>  	case OMP_CLAUSE_PRIVATE:
>  	case OMP_CLAUSE_SHARED:
>  	case OMP_CLAUSE_FIRSTPRIVATE:

Indentation.

> @@ -11056,6 +11097,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
>  	  WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
>  	  /* FALLTHRU */
>  
> +  case OMP_CLAUSE_INDEPENDENT:
>  	case OMP_CLAUSE_NOWAIT:
>  	case OMP_CLAUSE_ORDERED:
>  	case OMP_CLAUSE_DEFAULT:

Indentation.

> --- a/gcc/tree.def
> +++ b/gcc/tree.def
> @@ -1017,6 +1017,56 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
>  
>  DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
>  
> +/* #pragma acc kernels */
> +/* Operand 0: BODY
> +   Operand 1: CLAUSES   
> +*/
> +DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
> +
> +/* #pragma acc data */
> +/* Operand 0: BODY
> +   Operand 1: CLAUSES   
> +*/
> +DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
> +
> +/* #pragma acc host_data */
> +/* Operand 0: BODY
> +   Operand 1: CLAUSES   
> +*/
> +DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2)
> +
> +/* #pragma acc declare */
> +/* Operand 0: CLAUSES   
> +*/
> +DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
> +[...]
| [...]
| DEFTREECODE (OMP_PARALLEL, "omp_parallel", tcc_statement, 2)
| [...]
| DEFTREECODE (OMP_SINGLE, "omp_single", tcc_statement, 2)
| [...]
| DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
| [...]

If you look at tree.h:OMP_BODY and OMP_CLAUSES:

    #define OMP_BODY(NODE) \
      TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0)
    #define OMP_CLAUSES(NODE) \
      TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1)

..., you'll see that any codes from OACC_PARALLEL to OMP_CRITICAL are
expected to have a body as operand 0, and all from OACC_PARALLEL to
OMP_SINGLE are expected to have a list of clauses as operand 1.

This doesn't matter as long as all code paths are using the more
specialized OMP_*_BODY/OACC_*_BODY and OMP_*_CLAUSES/OACC_*_CLAUSES
accessor macros instead of the generic OMP_BODY and OMP_CLAUSES ones when
working with such OMP_* nodes, but to reduce confusion for readers of
this list/code, we should anyway pay attention, that is, distribute the
new OACC_* tree codes in tree.def to the appropriate places, to fit into
the existing scheme (as good as possible).  Please then also try to
maintain the same order of OACC_* codes in tree.h, tree.c, and so on:

> --- a/gcc/tree.h
> +++ b/gcc/tree.h
> @@ -1367,6 +1367,59 @@ extern void protected_set_expr_location (tree, location_t);
>  #define OMP_CLAUSE_DEFAULT_KIND(NODE) \
>    (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind)
>  
> +/* OpenACC directives and clause accessors. */
> +
> +#define OACC_KERNELS_BODY(NODE)      TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0)
> +#define OACC_KERNELS_CLAUSES(NODE)   TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1)

..., that is, move these (and the following ones) after the existing
OACC_PARALLEL_*, and so on, as appropriate.

> +#define OACC_DATA_BODY(NODE) \
> +            TREE_OPERAND (OACC_DATA_CHECK (NODE), 0)
> +
> +#define OACC_DATA_CLAUSES(NODE) \
> +            TREE_OPERAND (OACC_DATA_CHECK (NODE), 1)

No empty line between macros for the same tree code (OACC_DATA).

> +/* OpenACC clauses */

> +#define OMP_CLAUSE_[...]_EXPR(NODE) \

Please also try to sort these into the appropriate order (that is, as
defined in tree-core.h:omp_clause_code) -- if possible: I see that it is
not really consistent right now.  (This would also apply to the other
files where any OMP_*/OACC_* are used, for example the pretty printers,
but alas, if it's not in the right order now, then don't bother too much
about it.)


With these issues addressed, this patch is ready for commit to
gomp-4_0-branch.  Use your own judgement; if you feel confident, just
commit it, or otherwise post it again for a final review -- as you
prefer.


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]