Re: [gomp4] declare directive [3/5]

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi Jim!

I had mentioned that the Fortran front end changes cause regressions in a
few libgomp execution tests, if configured for Intel MIC (emulation)
offloading.  I have now located *where* this is coming from, but would
you please work on figuring out *why*?

Fortunately, you'll be able to work on the problem even without Intel MIC
(emulation) offloading configured: to reproduce, just look at the
difference in -fdump-tree-original without and with your patch applied.
You'll notice that clauses are getting "lost" from OpenMP target update
directives; for example, for
libgomp/testsuite/libgomp.fortran/declare-target-1.f90 I see:

    --- GOOD/declare-target-2.f90.003t.original     2015-06-16 18:16:07.472763339 +0200
    +++ ./declare-target-2.f90.003t.original        2015-06-16 19:28:22.706845250 +0200
    @@ -3,14 +3,14 @@
       extern integer(kind=4) var_x;
     
       var_x = 10;
    -  #pragma omp target update to(var_x)
    +  #pragma omp target update
       #pragma omp target
         {
           {
             var_x = var_x * 2;
           }
         }
    -  #pragma omp target update from(var_x)
    +  #pragma omp target update
       if (var_x != 20)
         {
           _gfortran_abort ();

(This is the only test case that I looked at, so far.)

I tracked this down to:

On Mon, 8 Jun 2015 10:04:11 -0500, James Norris <jnorris@codesourcery.com> wrote:
> --- a/gcc/fortran/trans-decl.c
> +++ b/gcc/fortran/trans-decl.c

> +static void
> +find_module_oacc_declare_clauses (gfc_symbol *sym)
> +{
> +  if (sym->attr.use_assoc)
> +    {
> +      gfc_omp_map_op map_op;
> +
> +      sym->attr.referenced = sym->attr.oacc_declare_create
> +			     | sym->attr.oacc_declare_copyin
> +			     | sym->attr.oacc_declare_deviceptr
> +			     | sym->attr.oacc_declare_device_resident;
> +
> +      if (sym->attr.oacc_declare_create)
> +	map_op = OMP_MAP_FORCE_ALLOC;
> +
> +      if (sym->attr.oacc_declare_copyin)
> +	map_op = OMP_MAP_FORCE_TO;
> +
> +      if (sym->attr.oacc_declare_deviceptr)
> +	map_op = OMP_MAP_FORCE_DEVICEPTR;
> +
> +      if (sym->attr.oacc_declare_device_resident)
> +	map_op = OMP_MAP_DEVICE_RESIDENT;
> +
> +      if (sym->attr.referenced)
> +	add_clause (sym, map_op);
> +    }
> +}

... this function apparently doing "something inappropriate".  It gets
(unconditionally) called from:

> +finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
>  {
> [...]
> +  gfc_traverse_ns (ns, find_module_oacc_declare_clauses);

... here, which in turn gets (unconditionally) called from:

> @@ -5946,8 +6237,7 @@ gfc_generate_function_code (gfc_namespace * ns)
>      add_argument_checking (&body, sym);
>  
>    /* Generate !$ACC DECLARE directive. */
> -  if (ns->oacc_declare)
> -    insert_oacc_declare (ns);
> +  finish_oacc_declare (ns, sym->attr.flavor);
>  
>    tmp = gfc_trans_code (ns->code);
>    gfc_add_expr_to_block (&body, tmp);

... here, and:

> --- a/gcc/fortran/trans-stmt.c
> +++ b/gcc/fortran/trans-stmt.c
> @@ -1588,8 +1588,7 @@ gfc_trans_block_construct (gfc_code* code)
>    code->exit_label = exit_label;
>  
>    /* Generate !$ACC DECLARE directive. */
> -  if (ns->oacc_declare)
> -    insert_oacc_declare (ns);
> +  finish_oacc_declare (ns, FL_UNKNOWN);
>  
>    gfc_add_expr_to_block (&body, gfc_trans_code (ns->code));
>    gfc_add_expr_to_block (&body, build1_v (LABEL_EXPR, exit_label));

... here.

Is that sufficient information for you to reproduce the problem?

As soon as you have a patch to bring back the lost clauses in the
-fdump-tree-original, I'll be happy to test it in my Intel MIC (emulated)
offloading build.


GrÃÃe,
 Thomas

Attachment: signature.asc
Description: PGP signature