Hi Jakub!

(I have not yet read any version of the OpenMP 4.1 standard.)

On Wed, 29 Apr 2015 13:14:06 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> --- gcc/tree.def.jj	2015-04-29 10:58:01.663745452 +0200
> +++ gcc/tree.def	2015-04-29 11:34:18.293302684 +0200

> +/* OpenMP - #pragma omp target enter data [clause1 ... clauseN]
> +   Operand 0: OMP_TARGET_ENTER_DATA_CLAUSES: List of clauses.  */
> +DEFTREECODE (OMP_TARGET_ENTER_DATA, "omp_target_enter_data", tcc_statement, 1)
> +
> +/* OpenMP - #pragma omp target exit data [clause1 ... clauseN]
> +   Operand 0: OMP_TARGET_EXIT_DATA_CLAUSES: List of clauses.  */
> +DEFTREECODE (OMP_TARGET_EXIT_DATA, "omp_target_exit_data", tcc_statement, 1)

;-) Heh, OpenMP catching up with OpenACC?

> --- gcc/c/c-parser.c.jj	2015-04-29 10:59:20.760504260 +0200
> +++ gcc/c/c-parser.c	2015-04-29 11:03:04.641991095 +0200

>     map ( variable-list )
>  
>     map-kind:
> -     alloc | to | from | tofrom  */
> +     alloc | to | from | tofrom
> +
> +   OpenMP 4.1:
> +   map-kind:
> +     alloc | to | from | tofrom | delete
> +
> +   map ( always [,] map-kind: variable-list ) */

"Funnily", OpenACC 2.5 is said to be dropping the "always" semantics that
OpenMP 4.1 is now adding...  That is, OpenACC 2.5's "copy" will then be
the same as "present_or_copy", and so on.

>  static tree
>  c_parser_omp_target_data (location_t loc, c_parser *parser)
>  {
> -  tree stmt = make_node (OMP_TARGET_DATA);
> -  TREE_TYPE (stmt) = void_type_node;
> -
> -  OMP_TARGET_DATA_CLAUSES (stmt)
> +  tree clauses
>      = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
>  				"#pragma omp target data");
> +  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
> +    {
> +      error_at (loc,
> +		"%<#pragma omp target data%> must contain at least one "
> +		"%<map%> clause");
> +      return NULL_TREE;
> +    }
> +
> +  tree stmt = make_node (OMP_TARGET_DATA);
> +  TREE_TYPE (stmt) = void_type_node;
> +  OMP_TARGET_DATA_CLAUSES (stmt) = clauses;

Even if it doesn't make a lot of sense, my reading of OpenMP 4.0 has been
that a target data construct without any map clauses is still valid.
(But I have not verified that now.)

> --- gcc/omp-low.c.jj	2015-04-29 10:58:01.489748182 +0200
> +++ gcc/omp-low.c	2015-04-29 11:03:04.646991017 +0200
> @@ -8994,6 +9017,11 @@ expand_omp_target (struct omp_region *re
>      case GF_OMP_TARGET_KIND_UPDATE:
>        start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
>        break;
> +    case GF_OMP_TARGET_KIND_ENTER_DATA:
> +    case GF_OMP_TARGET_KIND_EXIT_DATA:
> +      /* FIXME */
> +      start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
> +      break;

Actually, I've also wondered (but never verified) whether all of the
OpenACC enter, exit, and update constructs can be implemented via the
same libgomp API.

> @@ -11488,6 +11520,9 @@ lower_omp_target (gimple_stmt_iterator *
>  	      default:
>  		gcc_unreachable ();
>  	      }
> +	    /* FIXME: Temporary hack.  */
> +	    if (talign_shift == 3)
> +	      tkind &= ~GOMP_MAP_FLAG_FORCE;
>  	    gcc_checking_assert (tkind
>  				 < (HOST_WIDE_INT_C (1U) << talign_shift));
>  	    talign = ceil_log2 (talign);

Assuming you'll be changing the relevant functions' APIs, and assigning
new ABI versions, don't forget to also drop the unused offload_table
formal parameter (assuming that it remains unused).

> --- gcc/tree-pretty-print.c.jj	2015-04-29 10:58:01.663745452 +0200
> +++ gcc/tree-pretty-print.c	2015-04-29 11:03:04.648990986 +0200
> @@ -550,22 +560,22 @@ dump_omp_clause (pretty_printer *pp, tre
>  	  pp_string (pp, "tofrom");
>  	  break;
>  	case GOMP_MAP_FORCE_ALLOC:
> -	  pp_string (pp, "force_alloc");
> +	  pp_string (pp, "always,alloc");
>  	  break;
>  	case GOMP_MAP_FORCE_TO:
> -	  pp_string (pp, "force_to");
> +	  pp_string (pp, "always,to");
>  	  break;
>  	case GOMP_MAP_FORCE_FROM:
> -	  pp_string (pp, "force_from");
> +	  pp_string (pp, "always,from");
>  	  break;
>  	case GOMP_MAP_FORCE_TOFROM:
> -	  pp_string (pp, "force_tofrom");
> +	  pp_string (pp, "always,tofrom");
>  	  break;
>  	case GOMP_MAP_FORCE_PRESENT:
>  	  pp_string (pp, "force_present");
>  	  break;
>  	case GOMP_MAP_FORCE_DEALLOC:
> -	  pp_string (pp, "force_dealloc");
> +	  pp_string (pp, "always,delete");
>  	  break;
>  	case GOMP_MAP_FORCE_DEVICEPTR:
>  	  pp_string (pp, "force_deviceptr");

Makes some sense to me to use the same "always,*" syntax also for the
other "force_*" ones, given that these are artificial ("non-OpenACC")
descriptions anyway.


Grüße,
 Thomas