[Bug fortran/65419] New: incorrect sibcalls to libgomp functions

[Bug fortran/65419] New: incorrect sibcalls to libgomp functions

[cesar at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

            Bug ID: 65419
           Summary: incorrect sibcalls to libgomp functions
           Product: gcc
           Version: unknown
            Status: UNCONFIRMED
          Keywords: openacc
          Severity: normal
          Priority: P3
         Component: fortran
          Assignee: unassigned at gcc dot gnu.org
          Reporter: cesar at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org

I'm not sure if this problem is fortran specific, but I noticed that gfortran
is identifying calls to libgomp, specifically GOACC_data_end, as a sibcall if a
acc data region is the last statement inside the main function. That's a
problem because GOACC_data_end transfers data from the accelerator to the host
at the end of an acc data region. In the attached test case, the data in
question happens to be allocated on the stack. If you look at the epilogue of
the stack frame

    call    GOACC_data_start
    movq    %rbx, %rsi
    movq    %rsp, %rdi
    call    __field_summary_kernel_module_MOD_field_summary_kernel
    addq    $1296, %rsp
    .cfi_def_cfa_offset 16
    popq    %rbx
    .cfi_def_cfa_offset 8
# SUCC: EXIT [100.0%]  (ABNORMAL,SIBCALL)
    jmp    GOACC_data_end
    .cfi_endproc
.LFE1:

you'll see that %rsp is restored before GOACC_data_end is called. You should be
able to reproduce this bug with "-fopenacc -O2".

Did we specify the built-in declaration for GOACC_data_end wrong, or is this a
fortran specific quirk? I don't mind fixing it, but I'm not even sure what a
sibling call is. I think it has something to do with recursion.

Cesar

[Bug fortran/65419] incorrect sibcalls to libgomp functions

[cesar at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #2 from cesar at gcc dot gnu.org ---
Created attachment 35058
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=35058&action=edit
__builtin_GOACC_data_end

Sorry, I thought I included this test case. Anyway, the !$acc data regions is
associated with two libgomp calls, __builtin_GOACC_data_start and
__builtin_GOACC_data_end. Basically, as the control enters an acc data region,
all of the variables specified inside a data clause get transferred over to the
accelerator. At the end of the data region, GOACC_data_end transfers the data
back to the host as necessary. No arguments are passed to GOACC_data_end.
Instead the runtime maintains a stack of data regions because they are
permitted to nest, and GOACC_data_end decreases the reference count on all of
the variables pushed on that stack and transfers the data back to the host as
necessary.

Did we specify the builtin function for GOACC_data_end wrong, or is there an
inherent weakness in the sibling call analysis?

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

vries at gcc dot gnu.org changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |ASSIGNED
   Last reconfirmed|                            |2015-05-27
          Component|fortran                     |tree-optimization
     Ever confirmed|0                           |1

--- Comment #3 from vries at gcc dot gnu.org ---
Occurs on gomp-4_0-branch.

Minimal testcase in c:
...
$ cat src/libgomp/testsuite/libgomp.oacc-c-c++-common/data-end-sib-call.c
/* { dg-do run } */
/* { dg-options "-O2" } */

void
__attribute__((noinline,noclone))
f (void)
{
  int i;

#pragma acc data copyout (i)
  {

  }
}

int
main (void)
{
  f ();

  return 0;
}
...

The execution test passes on the host, but not on host_nonshm:
...
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-end-sib-call.c
-DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 execution test
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-end-sib-call.c
-DACC_DEVICE_TYPE_host_nonshm=1 -DACC_MEM_SHARED=0 execution test
...

The problem is in the declaration of GOACC_data_start:
...
DEF_GOACC_BUILTIN_FNSPEC (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
                          BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
                          ATTR_FNSPEC_DOT_DOT_r_r_r_NOTHROW_LIST,
                          ATTR_NOTHROW_LIST, "..rrr")
...

[ There's a problem with the matching.  The rs in "..rrr" were supposed to
match the PTR_PTR_PTR arguments. But that's not the case, since we need to add
a dot even for a void result. So the intended spec string was "...rrr".
AFAIU, this problem does not affect this PR. But I will review omp-builtins.def
for similar problems. ]

The 'r' means NOCLOBBER and NOESCAPE.

However, we take the address of 'i', assign it to .omp_data_arr.1 and pass that
to GOACC_data_start:
...
f ()
{
  int i;
  struct .omp_data_t.0 .omp_data_arr.1;
  static long unsigned int .omp_data_sizes.2[1] = {4};
  static short unsigned int .omp_data_kinds.3[1] = {642};

  <bb 2>:
  .omp_data_arr.1.i = &i;
  __builtin_GOACC_data_start (-1, 1, &.omp_data_arr.1, &.omp_data_sizes.2,
&.omp_data_kinds.3);
...

'i' escapes and is overwritten in __builtin_GOACC_data_start, so NOCLOBBER and
NOESCAPE do not hold.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #4 from vries at gcc dot gnu.org ---
When removing the fn spec from GOACC_data_start, we run into the problem that
this example doesn't get parallelized anymore:
...
#include <stdlib.h>

#define N (1024 * 512)
#define COUNTERTYPE unsigned int

int
main (void)
{
  unsigned int *__restrict a;
  unsigned int *__restrict b;
  unsigned int *__restrict c;

  a = (unsigned int *)malloc (N * sizeof (unsigned int));
  b = (unsigned int *)malloc (N * sizeof (unsigned int));
  c = (unsigned int *)malloc (N * sizeof (unsigned int));

  for (COUNTERTYPE i = 0; i < N; i++)
    a[i] = i * 2;

  for (COUNTERTYPE i = 0; i < N; i++)
    b[i] = i * 4;

#pragma acc data copyin (a[0:N], b[0:N]) copyout (c[0:N])
  {
#pragma acc kernels present (a[0:N], b[0:N], c[0:N])
    {
      for (COUNTERTYPE ii = 0; ii < N; ii++)
        c[ii] = a[ii] + b[ii];
    }
  }

  for (COUNTERTYPE i = 0; i < N; i++)
    if (c[i] != a[i] + b[i])
      abort ();

  free (a);
  free (b);
  free (c);

  return 0;
}
...

In this sequence, we take the address of a and pass it to GOACC_data_start:
...
  .omp_data_arr.18.a = &a;
  __builtin_GOACC_data_start (-1, 6, &.omp_data_arr.18, &.omp_data_sizes.19,
&.omp_data_kinds.20);
...

With the fnspec, we need to assume that a could be modified by
__builtin_GOACC_data_start. And that inhibits optimization.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #5 from vries at gcc dot gnu.org ---
This patch fixes it, but it's a bit of a hack:
...
diff --git a/gcc/tree-tailcall.c b/gcc/tree-tailcall.c
index 013972d..0cb73a7 100644
--- a/gcc/tree-tailcall.c
+++ b/gcc/tree-tailcall.c
@@ -499,6 +499,8 @@ find_tail_calls (basic_block bb, struct tailcall **ret)
     return;

   /* We found the call, check whether it is suitable.  */
+  if (gimple_call_builtin_p (stmt, BUILT_IN_GOACC_DATA_END))
+    return;
   tail_recursion = false;
   func = gimple_call_fndecl (call);
   if (func
...

A better solution is to add the omp_data_arr of the corresponding
GOACC_data_start as an argument to GOACC_data_end (with '.r' fnspec).

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[cesar at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #6 from cesar at gcc dot gnu.org ---
Isn't GOACC_parallel likely to have the same problem because hostaddrs may be
written to?

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #7 from vries at gcc dot gnu.org ---
Created attachment 35636
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=35636&action=edit
gcc/testsuite/c-c++-common/goacc/kernels-parallel-loop-data-enter-exit-2.c

(In reply to cesar from comment #6)
> Isn't GOACC_parallel likely to have the same problem because hostaddrs may
> be written to?

GOACC_parallel is not declared with a fnspec, so there's no tailcall problem:
...
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
  BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR,
  ATTR_NOTHROW_LIST)
...

OTOH, GOACC_parallel is an optimization barrier. So the attached testcase will
fail.  AFAICT, the only way this can be helped is by postponing expansion of
the parallel region until ompexpandssa, in other words do the same as for
kernels.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[jakub at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #8 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
IMNSHO it is much better to accurately describe the builtins to the aliasing
code etc. over adding ugly hacks like the tailcall one, or postponing expansion
till later etc.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #9 from vries at gcc dot gnu.org ---
(In reply to Jakub Jelinek from comment #8)
> IMNSHO it is much better to accurately describe the builtins to the aliasing
> code etc. over adding ugly hacks like the tailcall one, or postponing
> expansion till later etc.

We're discussing two issues now in the PR: GOACC_data_start/end and
GOACC_parallel. The issues are related because in both cases we're trying to
deal with a GOACC builtin acting as an optimization barrier for the kernels
region. The issues are also different because for GOACC_parallel there's an
associated code block, and for GOACC_data_start/end there isn't.

I'm working on fixing the GOACC_data_start/end issue by adding the omp_dara_arr
argument of GOACC_data_start to GOACC_data_end and adding fnspec ".r" to
GOACC_data_end. AFAIU, you agree with that approach.

The GOACC_parallel is more tricky. I don't see what kind of fnspec for
GOACC_parallel could fix this. The only other potential solution I see besides
postponing expansion of the parallel region is ipa_pta. But that's apparently
not production ready yet (PR46032 comment 19). Any other ideas?

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #10 from vries at gcc dot gnu.org ---
Created attachment 35639
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=35639&action=edit
tentative patch for GOACC_data_end

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[jakub at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #11 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
(In reply to vries from comment #10)
> Created attachment 35639 [details]
> tentative patch for GOACC_data_end

That is a bad idea, as void GOACC_data_end (void) is already exported out of
libgomp, so this is an ABI break.

If it is meant as a hack around not specifying the functions to aliasing oracle
properly, and not really needed after expansion to RTL, you could e.g. use an
internal function in GIMPLE and then lower it to normal function call that
doesn't take the argument.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #12 from vries at gcc dot gnu.org ---
Created attachment 35642
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=35642&action=edit
Updated tentative patch

(In reply to Jakub Jelinek from comment #11)
> (In reply to vries from comment #10)
> > Created attachment 35639 [details]
> > tentative patch for GOACC_data_end
> 
> That is a bad idea, as void GOACC_data_end (void) is already exported out of
> libgomp, so this is an ABI break.
> 
> If it is meant as a hack around not specifying the functions to aliasing
> oracle properly, and not really needed after expansion to RTL, you could
> e.g. use an internal function in GIMPLE and then lower it to normal function
> call that doesn't take the argument.

Updated tentative patch to use internal function.

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #13 from vries at gcc dot gnu.org ---
(In reply to vries from comment #9)
> The GOACC_parallel is more tricky. I don't see what kind of fnspec for
> GOACC_parallel could fix this. The only other potential solution I see
> besides postponing expansion of the parallel region is ipa_pta. But that's
> apparently not production ready yet (PR46032 comment 19). Any other ideas?

Filed as PR66324 - GOACC_parallel is optimization barrier

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

--- Comment #14 from vries at gcc dot gnu.org ---
Author: vries
Date: Thu May 28 14:18:19 2015
New Revision: 223832

URL: https://gcc.gnu.org/viewcvs?rev=223832&root=gcc&view=rev
Log:
Add IFN_GOACC_DATA_END_WITH_ARG

2015-05-28  Tom de Vries  <tom@codesourcery.com>

        PR tree-optimization/65419
        * cfgexpand.c (pass_data_expand): Add PROP_gimple_lompifn to
        properties_required field.
        * gimplify.c (gimplify_omp_workshare): Use IFN_GOACC_DATA_END_WITH_ARG
        instead of BUILT_IN_GOACC_DATA_END.  Clear PROP_gimple_lompifn in
        curr_properties.
        (gimplify_function_tree): Tentatively set PROP_gimple_lompifn in
        curr_properties.
        * internal-fn.c (expand_GOACC_DATA_END_WITH_ARG): New dummy function.
        * internal-fn.def (GOACC_DATA_END_WITH_ARG): New DEF_INTERNAL_FN.
        * omp-low.c (lower_omp_target): Set argument of
GOACC_DATA_END_WITH_ARG.
        (pass_data_late_lower_omp): New pass_data.
        (pass_late_lower_omp): New pass.
        (pass_late_lower_omp::gate, pass_late_lower_omp::execute)
        (make_pass_late_lower_omp): New function.
        * passes.def: Add pass_late_lower_omp.
        * tree-inline.c (expand_call_inline): Handle PROP_gimple_lompifn.
        * tree-pass.h (PROP_gimple_lompifn): Add define.

        * testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c: New test.

Added:
   
branches/gomp-4_0-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
Modified:
    branches/gomp-4_0-branch/gcc/ChangeLog.gomp
    branches/gomp-4_0-branch/gcc/cfgexpand.c
    branches/gomp-4_0-branch/gcc/gimplify.c
    branches/gomp-4_0-branch/gcc/internal-fn.c
    branches/gomp-4_0-branch/gcc/internal-fn.def
    branches/gomp-4_0-branch/gcc/omp-low.c
    branches/gomp-4_0-branch/gcc/passes.def
    branches/gomp-4_0-branch/gcc/tree-inline.c
    branches/gomp-4_0-branch/gcc/tree-pass.h
    branches/gomp-4_0-branch/libgomp/ChangeLog.gomp

[Bug tree-optimization/65419] incorrect sibcalls to libgomp functions

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=65419

vries at gcc dot gnu.org changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|ASSIGNED                    |RESOLVED
         Resolution|---                         |FIXED

--- Comment #16 from vries at gcc dot gnu.org ---
patch committed, marking resolved-fixed