From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-bugs-return-487383-listarch-gcc-bugs=gcc.gnu.org@gcc.gnu.org>
Received: (qmail 95262 invoked by alias); 27 May 2015 11:43:30 -0000
Mailing-List: contact gcc-bugs-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-bugs.gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-bugs/>
List-Post: <mailto:gcc-bugs@gcc.gnu.org>
List-Help: <mailto:gcc-bugs-help@gcc.gnu.org>
Sender: gcc-bugs-owner@gcc.gnu.org
Received: (qmail 95160 invoked by uid 48); 27 May 2015 11:43:24 -0000
From: "vries at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug tree-optimization/65419] incorrect sibcalls to libgomp functions
Date: Wed, 27 May 2015 11:43:00 -0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: tree-optimization
X-Bugzilla-Version: unknown
X-Bugzilla-Keywords: openacc
X-Bugzilla-Severity: normal
X-Bugzilla-Who: vries at gcc dot gnu.org
X-Bugzilla-Status: ASSIGNED
X-Bugzilla-Resolution:
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: vries at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags:
X-Bugzilla-Changed-Fields:
Message-ID: <bug-65419-4-ixEoOvMXGv@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-65419-4@http.gcc.gnu.org/bugzilla/>
References: <bug-65419-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: 7bit
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
X-SW-Source: 2015-05/txt/msg02223.txt.bz2

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.