Re: [PATCH] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct

[Jakub Jelinek <jakub@redhat.com>]

On Wed, Nov 24, 2021 at 06:08:02PM +0100, Marcel Vollweiler wrote:
> +	case OMP_CLAUSE_HAS_DEVICE_ADDR:
> +	  t = OMP_CLAUSE_DECL (c);
> +	  if (TREE_CODE (t) == TREE_LIST)
> +	    {
> +	      if (handle_omp_array_sections (c, ort))
> +		remove = true;
> +	      else
> +		{
> +		  t = OMP_CLAUSE_DECL (c);
> +		  while (TREE_CODE (t) == ARRAY_REF)
> +		    t = TREE_OPERAND (t, 0);
> +		}
> +	    }
> +	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
> +	    bitmap_set_bit (&is_on_device_head, DECL_UID (t));

Why the OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR check?
There is no goto into this block nor fallthru into it, and
handle_omp_array_sections better shouldn't change OMP_CLAUSE_CODE.

>  	  goto check_dup_generic;
>  
> +	case OMP_CLAUSE_HAS_DEVICE_ADDR:
> +	  t = OMP_CLAUSE_DECL (c);
> +	  if (TREE_CODE (t) == TREE_LIST)
> +	    if (handle_omp_array_sections (c, ort))
> +	      remove = true;
> +	    else
> +	      {
> +		t = OMP_CLAUSE_DECL (c);
> +		while (TREE_CODE (t) == ARRAY_REF)
> +		  t = TREE_OPERAND (t, 0);
> +	      }
> +	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
> +	    bitmap_set_bit (&is_on_device_head, DECL_UID (t));

Likewise.

> +	  if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
> +	    cxx_mark_addressable (t);
> +	  goto check_dup_generic_t;
> +
>  	case OMP_CLAUSE_USE_DEVICE_ADDR:
>  	  field_ok = true;
>  	  t = OMP_CLAUSE_DECL (c);

> --- a/gcc/fortran/gfortran.h
> +++ b/gcc/fortran/gfortran.h
> @@ -1391,7 +1391,8 @@ enum
>    OMP_LIST_USE_DEVICE_PTR,
>    OMP_LIST_USE_DEVICE_ADDR,
>    OMP_LIST_NONTEMPORAL,
> -  OMP_LIST_NUM
> +  OMP_LIST_HAS_DEVICE_ADDR,
> +  OMP_LIST_NUM  /* must be the last  */

Capital M and . at the end.

> @@ -2077,6 +2078,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>  	    }
>  	  break;
>  	case 'h':
> +	  if ((mask & OMP_CLAUSE_HAS_DEVICE_ADDR)
> +	      && gfc_match_omp_variable_list
> +		   ("has_device_addr (",
> +		    &c->lists[OMP_LIST_HAS_DEVICE_ADDR], false, NULL, NULL,
> +							 true) == MATCH_YES)

Formatting, true should be IMO below &c->lists.

> +	    continue;
>  	  if ((mask & OMP_CLAUSE_HINT)
>  	      && (m = gfc_match_dupl_check (!c->hint, "hint", true, &c->hint))
>  		 != MATCH_NO)
> @@ -2850,7 +2857,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>  	  if ((mask & OMP_CLAUSE_USE_DEVICE_ADDR)
>  	      && gfc_match_omp_variable_list
>  		   ("use_device_addr (",
> -		    &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES)
> +		    &c->lists[OMP_LIST_USE_DEVICE_ADDR], false, NULL, NULL,
> +							 true) == MATCH_YES)

Likewise.

> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -1910,7 +1910,17 @@ gfc_trans_omp_variable_list (enum omp_clause_code code,
>  	tree t = gfc_trans_omp_variable (namelist->sym, declare_simd);
>  	if (t != error_mark_node)
>  	  {
> -	    tree node = build_omp_clause (input_location, code);
> +	    tree node;
> +	    /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the
> +	       descriptor such that the bounds are available; its data component
> +	       is unmodified; it is handled as device address inside target. */
> +	    if (code == OMP_CLAUSE_HAS_DEVICE_ADDR
> +		&& (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (t))
> +		    || (POINTER_TYPE_P (TREE_TYPE (t))
> +			&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (t))))))
> +	      node = build_omp_clause (input_location, OMP_CLAUSE_FIRSTPRIVATE);

Not sure about the above,

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -10024,6 +10024,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	  flags = GOVD_EXPLICIT;
>  	  goto do_add;
>  
> +	case OMP_CLAUSE_HAS_DEVICE_ADDR:
> +	  decl = OMP_CLAUSE_DECL (c);
> +	  if (TREE_CODE (decl) == ARRAY_REF)
> +	    {
> +	      flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
> +	      while (TREE_CODE (decl) == ARRAY_REF)
> +		decl = TREE_OPERAND (decl, 0);
> +	      goto do_add_decl;

but this looks weird.
If decl after stripping the ARRAY_REFs is a var with pointer type, sure,
firstprivatizing it is the way to go.
But it can be also a variable with ARRAY_TYPE, can't it?  Something like:
  int a[64];
  #pragma omp target data map(a) use_device_addr(a)
  {
    #pragma omp target has_device_addr(a[3:16])
    a[3] = 1;
  }
and in this case firstprivatization of a looks wrong.  use_device_addr
should replace (but only at omp-low.c time I think) a used in the block
with the remapped a (i.e. *device_address_of_a).
Or perhaps it could be a non-static data member with array type
inside of a C++ method.

> +	case OMP_CLAUSE_HAS_DEVICE_ADDR:
> +	  decl = OMP_CLAUSE_DECL (c);
> +	  if (TREE_CODE (decl) == ARRAY_REF)
> +	    while (TREE_CODE (decl) == ARRAY_REF)
> +	      decl = TREE_OPERAND (decl, 0);

Isn't this equivalent to just the while loop without the if?

	Jakub