On Mon, 2 Nov 2015 18:33:39 +0000
Julian Brown <julian@codesourcery.com> wrote:

> On Mon, 26 Oct 2015 19:34:22 +0100
> Jakub Jelinek <jakub@redhat.com> wrote:
> 
> > Your use_device sounds very similar to use_device_ptr clause in
> > OpenMP, which is allowed on #pragma omp target data construct and is
> > implemented quite a bit differently from this; it is unclear if the
> > OpenACC standard requires this kind of implementation, or you just
> > chose to implement it this way.  In particular, the GOMP_target_data
> > call puts the variables mentioned in the use_device_ptr clauses into
> > the mapping structures (similarly how map clause appears) and the
> > corresponding vars are privatized within the target data region
> > (which is a host region, basically a fancy { } braces), where the
> > private variables contain the offloading device's pointers.  
> 
> As the author of the original patch, I have to say using the mapping
> structures seems like a far better approach, but I've hit some trouble
> with the details of adapting OpenACC to use that method.

Here's a version of the patch which (hopefully) brings OpenACC on par
with OpenMP with respect to use_device/use_device_ptr variables. The
implementation is essentially the same now for OpenACC as for OpenMP
(i.e. using mapping structures): so for now, only array or pointer
variables can be used as use_device variables. The included tests have
been adjusted accordingly.

One awkward part of the implementation concerns nesting offloaded
regions within host_data regions:

#define N 1024

int main (int argc, char* argv[])
{
  int x[N];

#pragma acc data copyin (x[0:N])
  {
    int *xp;
#pragma acc host_data use_device (x)
    {
      [...]
#pragma acc parallel present (x) copyout (xp)
      {
        xp = x;
      }
    }

    assert (xp == acc_deviceptr (x));
  }

  return 0;
}

I think the meaning of 'x' as seen within the clauses of the parallel
directive should be the *host* version of x, not the mapped target
address (I've asked on the OpenACC technical mailing list to clarify
this point, but no reply as yet). The changes to
{maybe_,}lookup_decl_in_outer_ctx "skip over" host_data contexts when
called from lower_omp_target. There's probably an analogous case for
OpenMP, but I've not tried to handle that.

No regressions for libgomp tests, and the new tests pass. OK for trunk?

Thanks,

Julian

ChangeLog

    Julian Brown  <julian@codesourcery.com>
    Cesar Philippidis  <cesar@codesourcery.com>
    James Norris  <James_Norris@mentor.com>

    gcc/
    * c-family/c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
    * c-family/c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.
    (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_USE_DEVICE.
    * c/c-parser.c (c_parser_omp_clause_name): Add use_device support.
    (c_parser_oacc_clause_use_device): New function.
    (c_parser_oacc_all_clauses): Add use_device support.
    (OACC_HOST_DATA_CLAUSE_MASK): New macro.
    (c_parser_oacc_host_data): New function.
    (c_parser_omp_construct): Add host_data support.
    * c/c-tree.h (c_finish_oacc_host_data): Add prototype.
    * c/c-typeck.c (c_finish_oacc_host_data): New function.
    (c_finish_omp_clauses): Add use_device support.
    * cp/cp-tree.h (finish_oacc_host_data): Add prototype.
    * cp/parser.c (cp_parser_omp_clause_name): Add use_device support.
    (cp_parser_oacc_all_clauses): Add use_device support.
    (OACC_HOST_DATA_CLAUSE_MASK): New macro.
    (cp_parser_oacc_host_data): New function.
    (cp_parser_omp_construct): Add host_data support.
    (cp_parser_pragma): Add host_data support.
    * cp/semantics.c (finish_omp_clauses): Add use_device support.
    (finish_oacc_host_data): New function.
    * gimple-pretty-print.c (dump_gimple_omp_target): Add host_data
    support.
    * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_HOST_DATA.
    (is_gimple_omp_oacc): Add support for above.
    * gimplify.c (gimplify_scan_omp_clauses): Add host_data, use_device
    support.
    (gimplify_omp_workshare): Add host_data support.
    (gimplify_expr): Likewise.
    * omp-builtins.def (BUILT_IN_GOACC_HOST_DATA): New.
    * omp-low.c (lookup_decl_in_outer_ctx)
    (maybe_lookup_decl_in_outer_ctx): Add optional argument to skip
    host_data regions.
    (scan_sharing_clauses): Support use_device.
    (check_omp_nesting_restrictions): Support host_data.
    (expand_omp_target): Support host_data.
    (lower_omp_target): Skip over outer host_data regions when looking
    up decls. Support use_device.
    (make_gimple_omp_edges): Support host_data.
    * tree-nested.c (convert_nonlocal_omp_clauses): Add use_device
    clause.

    libgomp/
    * oacc-parallel.c (GOACC_host_data): New function.
    * libgomp.map (GOACC_host_data): Add to GOACC_2.0.1.
    * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/host_data-2.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/host_data-3.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/host_data-4.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/host_data-5.c: New test.