On Thu, 19 Nov 2015 16:57:23 +0100
Jakub Jelinek <jakub@redhat.com> wrote:

> If it is unclear, I think disallowing acc {parallel,kernels} inside of
> acc host_data might be too big hammer, but perhaps just erroring out
> or warning during gimplification that if you (explicitly or
> implicitly) try to map a var that is in use_device clause in some
> outer context, it is either wrong, unsupported or will not do what
> users think?

I think we can only assume that trying to map a variable declared in
a surrounding use_device clause is undefined behaviour. I haven't had
any response to my questions about host_data & deviceptr on the OpenACC
list.

> > #pragma acc host_data use_device(x)
> > {
> >   target_primitive(x);
> >   #pragma acc parallel deviceptr(x)
> >   {
> >     ...
> >   }
> > }
> 
> Is deviceptr as above meant to work?  That is the OpenACC counterpart
> of is_device_ptr, right?  If yes, then I'd suggest just warning if you
> try to implicitly or explicitly map something use_device in outer
> contexts, and just make sure you don't ICE on the cases where you
> warn. If the standard does not say what it means, then it is
> unspecified behavior...

A problem with deviceptr, unlike is_device_ptr, is that it turns out to
be defined only to work with pointers, not arrays (OpenACC 2.0a
2.6.5.2), and there are no rules describing the latter decaying to the
former. So at least if 'x' is an array, it appears the answer is "no".

So, the attached patch disallows (via raising an error):

* Variables being declared in explicit mapping clauses that are
  declared in enclosing host_data regions.

* Variables being implicitly used (mapped) in offloaded regions that
  are declared in enclosing host_data regions.

It's otherwise equivalent to the previously-posted version, but without
the hacks to {maybe_,}lookup_decl_in_outer_ctx. I added checks for the
above conditions during gimplification, which seemed to be about the
same phase that other similar kinds of errors are diagnosed.

Tests look OK (libgomp/gcc/g++/libstdc++), and the new ones pass.

OK for mainline?

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 (omp_region_type): Add ORT_ACC_HOST_DATA.
    (omp_notice_variable): Diagnose undefined implicit uses of
    use_device variables in offloaded regions.
    (gimplify_scan_omp_clauses): Add host_data, use_device
    support. Diagnose undefined mapping of use_device variables in
    OpenACC clauses.
    (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.
    * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.