diff mbox

[gomp4] fix implicit data clause linker error

Message ID 068721cc-1150-385e-7c34-229d50488ea2@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Dec. 6, 2016, 5:20 p.m. UTC 
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
diff mbox

Patch

2016-12-06  Cesar Philippidis  <cesar@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" } }