gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if...

	* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
	GOVD_EXPLICIT flags.
	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
	call install_var_field with mask 11 instead of 3.
	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
gcc/cp/
	* semantics.c (finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
libgomp/
	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
	perform the lookup in the first loop only if !not_found_cnt, otherwise
	perform lookups for it in the second loop guarded with
	if (not_found_cnt || has_firstprivate).
	* testsuite/libgomp.c/target-37.c: New test.
	* testsuite/libgomp.c++/target-22.C: New test.

From-SVN: r274206

1 files changed, 71 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
new file mode 100644
index 0000000..b3cc6a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-37.c
@@ -0,0 +1,71 @@
+extern void abort (void);
+struct S { int e, f; };
+
+void
+foo (int n)
+{
+  int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
+  struct S d = { 5, 6 };
+  int *p = a + 1, i, err;
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data use_device_ptr(p) map(from:err) map(to:a)
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < 4; i++)
+    a[i] = 23 + i;
+  #pragma omp target data map(to:a) use_device_addr(a) map(from:err)
+  #pragma omp target is_device_ptr(a) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (a[i] != 23 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr(b) map(from:err) map(to:b)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:c) use_device_addr(c) map(from:err)
+  {
+    int *q = &c;
+    #pragma omp target is_device_ptr(q) map(from:err)
+    {
+      err = *q != 4;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr(d) map(to:d) map(from:err)
+  {
+    struct S *r = &d;
+    #pragma omp target is_device_ptr(r) map(from:err)
+    {
+      err = r->e != 5 || r->f != 6;
+    }
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}