tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description.

	* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
	OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
	for OMP_CLAUSE_USE_DEVICE_ADDR clause.
	(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
	Likewise.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
	Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
	clause with array or reference to array types, no matter what type
	except for reference it has.
gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
	equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
	enumeration value.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
	(c_parser_omp_clause_use_device_addr): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer type rather than pointer or array type.
	Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer or reference to pointer type rather than
	pointer or array or reference to pointer or array type. Handle
	OMP_CLAUSE_USE_DEVICE_ADDR.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/testsuite/
	* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
	instead of use_device_ptr clause where required by OpenMP 5.0, add
	further tests for both use_device_ptr and use_device_addr clauses.
libgomp/
	* testsuite/libgomp.c/target-18.c (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.
	* testsuite/libgomp.c++/target-9.C (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.  Add t and u arguments.
	(main): Adjust caller.

From-SVN: r274159

3 files changed, 77 insertions, 8 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0331c51..01440f7 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,15 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+	* testsuite/libgomp.c/target-18.c (struct S): New type.
+	(foo): Use use_device_addr clause instead of use_device_ptr clause
+	where required by OpenMP 5.0, add further tests for both use_device_ptr
+	and use_device_addr clauses.
+	* testsuite/libgomp.c++/target-9.C (struct S): New type.
+	(foo): Use use_device_addr clause instead of use_device_ptr clause
+	where required by OpenMP 5.0, add further tests for both use_device_ptr
+	and use_device_addr clauses.  Add t and u arguments.
+	(main): Adjust caller.
+
 2019-08-06  Jakub Jelinek  <jakub@redhat.com>
 
 	* testsuite/libgomp.c++/loop-13.C: New test.
diff --git a/libgomp/testsuite/libgomp.c++/target-9.C b/libgomp/testsuite/libgomp.c++/target-9.C
index a5d171b..83a61cf 100644
--- a/libgomp/testsuite/libgomp.c++/target-9.C
+++ b/libgomp/testsuite/libgomp.c++/target-9.C
@@ -1,10 +1,13 @@
 extern "C" void abort (void);
+struct S { int e, f; };
 
 void
-foo (int *&p, int (&s)[5], int n)
+foo (int *&p, int (&s)[5], int &t, S &u, int n)
 {
   int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
   int *r = a + 1, *q = p - 1, i, err;
+  int v = 27;
+  S w = { 28, 29 };
   for (i = 0; i < n; i++)
     b[i] = 9 + i;
   #pragma omp target data map(to:a)
@@ -30,7 +33,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:b)
-  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target data use_device_addr(b) map(from:err)
   #pragma omp target is_device_ptr(b) private(i) map(from:err)
   {
     err = 0;
@@ -41,7 +44,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:c)
-  #pragma omp target data use_device_ptr(c) map(from:err)
+  #pragma omp target data use_device_addr(c) map(from:err)
   #pragma omp target is_device_ptr(c) private(i) map(from:err)
   {
     err = 0;
@@ -52,7 +55,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:s[:5])
-  #pragma omp target data use_device_ptr(s) map(from:err)
+  #pragma omp target data use_device_addr(s) map(from:err)
   #pragma omp target is_device_ptr(s) private(i) map(from:err)
   {
     err = 0;
@@ -62,6 +65,34 @@ foo (int *&p, int (&s)[5], int n)
   }
   if (err)
     abort ();
+  #pragma omp target data map(to: v) map(to:u)
+  #pragma omp target data use_device_addr (v) use_device_addr (u) map(from:err)
+  {
+    int *z = &v;
+    S *x = &u;
+    #pragma omp target is_device_ptr (z, x) map(from:err)
+    {
+      err = 0;
+      if (*z != 27 || x->e != 25 || x->f != 26)
+	err = 1;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to: t, w)
+  #pragma omp target data use_device_addr (t, w) map(from:err)
+  {
+    int *z = &t;
+    S *x = &w;
+    #pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
+    {
+      err = 0;
+      if (*z != 24 || x->e != 28 || x->f != 29)
+	err = 1;
+    }
+  }
+  if (err)
+    abort ();
 }
 
 int
@@ -69,5 +100,7 @@ main ()
 {
   int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
   int *p = a + 1;
-  foo (p, b, 9);
+  int t = 24;
+  S u = { 25, 26 };
+  foo (p, b, t, u, 9);
 }
diff --git a/libgomp/testsuite/libgomp.c/target-18.c b/libgomp/testsuite/libgomp.c/target-18.c
index cbacaf6..dd511fb 100644
--- a/libgomp/testsuite/libgomp.c/target-18.c
+++ b/libgomp/testsuite/libgomp.c/target-18.c
@@ -1,9 +1,11 @@
 extern void abort (void);
+struct S { int e, f; };
 
 void
 foo (int n)
 {
-  int a[4] = { 0, 1, 2, 3 }, b[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;
@@ -21,7 +23,7 @@ foo (int n)
   for (i = 0; i < 4; i++)
     a[i] = 23 + i;
   #pragma omp target data map(to:a)
-  #pragma omp target data use_device_ptr(a) map(from:err)
+  #pragma omp target data use_device_addr(a) map(from:err)
   #pragma omp target is_device_ptr(a) private(i) map(from:err)
   {
     err = 0;
@@ -32,7 +34,7 @@ foo (int n)
   if (err)
     abort ();
   #pragma omp target data map(to:b)
-  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target data use_device_addr(b) map(from:err)
   #pragma omp target is_device_ptr(b) private(i) map(from:err)
   {
     err = 0;
@@ -42,6 +44,28 @@ foo (int n)
   }
   if (err)
     abort ();
+  #pragma omp target data map(to:c)
+  #pragma omp target data 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 map(to:d)
+  #pragma omp target data use_device_addr(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