The patch I introduced here <https://gcc.gnu.org/ml/gcc-patches/2016-08/msg01293.html> to fix PR70828 in gomp-4_0-branch has a bug involving the creation of new implicit data clauses for variables inside ACC PARALLEL or KERNELS regions that are present inside an enclosing ACC DATA region. Specifically, as gimplifiy_adjust_omp_clauses_1 introduced new present data clauses for those variables, it neglected to link those clauses with the existing clauses. This is bad in any case, but for data clauses involving static or global variables, this would result in an LTO error because those global variable definitions aren't included inside the offloaded LTO partitions.
I've applied this fix to gomp-4_0-branch. Cesar
2016-12-06 Cesar Philippidis <ce...@codesourcery.com> gcc/ * gimplify.c (gimplify_adjust_omp_clauses_1): Link ACC new clauses with the old ones. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 9649fae..00b1b23 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7992,6 +7992,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) gimplify_expr (&OMP_CLAUSE_SIZE (nc), pre_p, NULL, is_gimple_val, fb_rvalue); gimplify_omp_ctxp = ctx; + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); OMP_CLAUSE_CHAIN (clause) = nc; } else if (DECL_SIZE (decl) diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c new file mode 100644 index 0000000..c2f33ff --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c @@ -0,0 +1,24 @@ +/* Ensure that the gimplifier does not remove any existing clauses as + it inserts new implicit data clauses. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 +static int a[N], b[N]; + +int main(int argc, char *argv[]) +{ + int i; + +#pragma acc data copyin(a[0:N]) copyout (b[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + b[i] = a[i]; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times "omp target oacc_data map.force_from:b.0. .len: 400.. map.force_to:a.0. .len: 400.." 1 "gimple" } } +// { dg-final { scan-tree-dump-times "omp target oacc_parallel map.force_present:b.0. .len: 400.. map.firstprivate:b .pointer assign, bias: 0.. map.force_present:a.0. .len: 400.. map.firstprivate:a .pointer assign, bias: 0.." 1 "gimple" } }