Hi Jakub,

>> +    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.

Good point. Removed.

>
>>        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.

Removed.

>
>> +      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.

Changed.

>
>> @@ -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.

Corrected the formatting.

>
>> +        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.

Corrected.

>
>> --- 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,

This is needed for allocatable arrays and array pointers to ensure that
not only the (array) data is (already) present on the device but also
the array descriptor. Otherwise the test cases
target-has-device-addr-2.f90, target-has-device-addr-3.f90 (because of
variable "c") and target-has-device-addr-4.f90 (also because of variable
"c") won't work.

>
>> --- 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.
I removed GOVD_FIRSTPRIVATE from the OMP_CLAUSE_HAS_DEVICE_ADDR case and
simplified the above code to:

        case OMP_CLAUSE_HAS_DEVICE_ADDR:
          decl = OMP_CLAUSE_DECL (c);
          while (TREE_CODE (decl) == INDIRECT_REF
                 || TREE_CODE (decl) == ARRAY_REF)
            decl = TREE_OPERAND (decl, 0);
          flags = GOVD_EXPLICIT;
          goto do_add_decl;

The reason is that I don't see a case where privatizing is useful,
because whatever type a list item has, we claim with has_device_addr
that the list item is already on the device. So we shouldn't create a
duplicate with a new device address for any possible list item.

>
>> +    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?

Yes. I removed the if statement.

Moreover, I added testing for references to array sections (without
enclosing target data region) and for references to pointers. To tackle
also those cases further small code revisions were necessary.

The patch was tested with x86-64 and powerpc configs with offoading
(nvptx and gcn) without regressions.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955