OpenMP: Enable has_device_addr clause for 'dispatch' in C/C++

The 'has_device_addr' of 'dispatch' has to be seen in conjunction with the
'need_device_addr' modifier to the 'adjust_args' clause of 'declare variant'.
As the latter has not yet been implemented, 'has_device_addr' has no real
effect. However, to prepare for 'need_device_addr' and as service to the user:

For C, where 'need_device_addr' is not permitted (contrary to C++ and Fortran),
a note is output when then the user tries to use it (alongside the existing
error that either 'nothing' or 'need_device_ptr' was expected).

And, on the ME side, is is lightly handled by diagnosing when - for the
same argument - there is a mismatch between the variant's adjust_args
'need_device_ptr' modifier and dispatch having an 'has_device_addr' clause
(or likewise for need_device_addr with is_device_ptr) as, according to the
spec, those are completely separate.
Thus, 'dispatch' will still do the host to device pointer conversion for
a 'need_device_ptr' argument, even if it appeared in a 'has_device_addr'
clause.

gcc/c/ChangeLog:

	* c-parser.cc (OMP_DISPATCH_CLAUSE_MASK): Add has_device_addr clause.
	(c_finish_omp_declare_variant): Add an 'inform' telling the user that
	'need_device_addr' is invalid for C.

gcc/cp/ChangeLog:

	* parser.cc (OMP_DISPATCH_CLAUSE_MASK): Add has_device_addr clause.

gcc/ChangeLog:

	* gimplify.cc (gimplify_call_expr): When handling OpenMP's dispatch,
	add diagnostic when there is a ptr vs. addr mismatch between
	need_device_{addr,ptr} and {is,has}_device_{ptr,addr}, respectively.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/adjust-args-3.c: New test.
	* gcc.dg/gomp/adjust-args-2.c: New test.

5 files changed, 152 insertions, 18 deletions

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 4ec0ee8..d023580 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -25271,6 +25271,7 @@ c_parser_omp_dispatch_body (c_parser *parser)
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)                           \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS)                       \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT)                        \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)                  \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP)                          \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)                    \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
@@ -26963,6 +26964,9 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
 		{
 		  error_at (c_parser_peek_token (parser)->location,
 			    "expected %<nothing%> or %<need_device_ptr%>");
+		  if (strcmp (p, "need_device_addr") == 0)
+		    inform (c_parser_peek_token (parser)->location,
+			    "%<need_device_addr%> is not valid for C");
 		  goto fail;
 		}
 	    }
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 8a3472a..15a5253 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -49924,6 +49924,7 @@ cp_parser_omp_dispatch_body (cp_parser *parser)
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)                           \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS)                       \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT)                        \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)                  \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP)                          \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)                    \
    | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 8495c45..dd0d992 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -4124,27 +4124,39 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
 			arg_types = TREE_CHAIN (arg_types);
 
 		      bool need_device_ptr = false;
-		      for (tree arg
-			   = TREE_PURPOSE (TREE_VALUE (adjust_args_list));
-			   arg != NULL; arg = TREE_CHAIN (arg))
-			{
-			  if (TREE_VALUE (arg)
-			      && TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST
-			      && wi::eq_p (i, wi::to_wide (TREE_VALUE (arg))))
-			    {
-			      need_device_ptr = true;
-			      break;
-			    }
-			}
+		      bool need_device_addr = false;
+		      for (int need_addr = 0; need_addr <= 1; need_addr++)
+			for (tree arg = need_addr
+					? TREE_VALUE (TREE_VALUE (
+					    adjust_args_list))
+					: TREE_PURPOSE (TREE_VALUE (
+					    adjust_args_list));
+			     arg != NULL; arg = TREE_CHAIN (arg))
+			  {
+			    if (TREE_VALUE (arg)
+				&& TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST
+				&& wi::eq_p (i, wi::to_wide (TREE_VALUE (arg))))
+			      {
+				if (need_addr)
+				  need_device_addr = true;
+				else
+				  need_device_ptr = true;
+				break;
+			      }
+			  }
 
-		      if (need_device_ptr)
+		      if (need_device_ptr || need_device_addr)
 			{
 			  bool is_device_ptr = false;
+			  bool has_device_addr = false;
+
 			  for (tree c = gimplify_omp_ctxp->clauses; c;
 			       c = TREE_CHAIN (c))
 			    {
-			      if (OMP_CLAUSE_CODE (c)
-				  == OMP_CLAUSE_IS_DEVICE_PTR)
+			      if ((OMP_CLAUSE_CODE (c)
+				   == OMP_CLAUSE_IS_DEVICE_PTR)
+				  || (OMP_CLAUSE_CODE (c)
+				      == OMP_CLAUSE_HAS_DEVICE_ADDR))
 				{
 				  tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c));
 				  tree decl2
@@ -4155,15 +4167,42 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
 				      || TREE_CODE (decl2) == PARM_DECL)
 				    {
 				      decl2 = DECL_NAME (decl2);
-				      if (decl1 == decl2)
-					is_device_ptr = true;
+				      if (decl1 == decl2
+					  && (OMP_CLAUSE_CODE (c)
+					      == OMP_CLAUSE_IS_DEVICE_PTR))
+					{
+					  if (need_device_addr)
+					    warning_at (
+					      OMP_CLAUSE_LOCATION (c),
+					      OPT_Wopenmp,
+					      "%<is_device_ptr%> for %qD does"
+					      " not imply %<has_device_addr%> "
+					      "required for "
+					      "%<need_device_addr%>",
+					       OMP_CLAUSE_DECL (c));
+					  is_device_ptr = true;
+					}
+				      else if (decl1 == decl2)
+					{
+					  if (need_device_ptr)
+					    warning_at (
+					      OMP_CLAUSE_LOCATION (c),
+					      OPT_Wopenmp,
+					      "%<has_device_addr%> for %qD does"
+					      " not imply %<is_device_ptr%> "
+					      "required for "
+					      "%<need_device_ptr%>",
+					      OMP_CLAUSE_DECL (c));
+					  has_device_addr = true;
+					}
 				    }
 				}
 			      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE)
 				device_num = OMP_CLAUSE_OPERAND (c, 0);
 			    }
 
-			  if (!is_device_ptr)
+			  if ((need_device_ptr && !is_device_ptr)
+			      || (need_device_addr && !has_device_addr))
 			    {
 			      if (device_num == NULL_TREE)
 				{
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c
new file mode 100644
index 0000000..f62272c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c
@@ -0,0 +1,85 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+// Do diagnostic check / dump check only;
+// Note: this test should work as run-test as well.
+
+#if 0
+  #include <omp.h>
+#else
+  #ifdef __cplusplus
+  extern "C" {
+  #endif
+    extern int omp_get_default_device ();
+    extern int omp_get_num_devices ();
+  #ifdef __cplusplus
+  }
+  #endif
+#endif
+
+
+void f(int *x, int *y);
+#pragma omp declare variant(f) adjust_args(need_device_ptr: x, y) match(construct={dispatch})
+void g(int *x, int *y);
+
+void
+sub (int *a, int *b)
+{
+  // The has_device_addr is a bit questionable as the caller is not actually
+  // passing a device address - but we cannot pass one because of the
+  // following:
+  //
+  // As for 'b' need_device_ptr has been specified and 'b' is not
+  // in the semantic requirement set 'is_device_ptr' (and only in 'has_device_addr')
+  // "the argument is converted in the same manner that a use_device_ptr clause
+  //  on a target_data construct converts its pointer"
+  #pragma omp dispatch is_device_ptr(a), has_device_addr(b)  /* { dg-warning "'has_device_addr' for 'b' does not imply 'is_device_ptr' required for 'need_device_ptr' \\\[-Wopenmp\\\]" } */
+    g(a, b);
+}
+
+void
+f(int *from, int *to)
+{
+  static int cnt = 0;
+  cnt++;
+  if (cnt >= 3)
+    {
+      if (omp_get_default_device () != -1
+          && omp_get_default_device () < omp_get_num_devices ())
+        {
+	  // On offload device but not mapped
+	  if (from != (void *)0L) // Not mapped
+	    __builtin_abort ();
+        }
+      else if (from[0] != 5)
+        __builtin_abort ();
+      return;
+    }
+  #pragma omp target is_device_ptr(from, to)
+  {
+    to[0] = from[0] * 10;
+    to[1] = from[1] * 10;
+  }
+}
+
+int
+main ()
+{
+  int A[2], B[2] = {123, 456}, C[1] = {5};
+  int *p = A;
+  #pragma omp target enter data map(A, B)
+
+  /* Note: We don't add  'use_device_addr(B)' here;
+     if we do, it will fail with an illegal memory access (why?).  */
+  #pragma omp target data use_device_ptr(p)
+    {
+      sub(p, B);
+      sub(C, B); /* C is not mapped -> 'from' ptr == NULL  */
+    }
+
+  #pragma omp target exit data map(A, B)
+}
+
+// { dg-final { scan-tree-dump-times "#pragma omp dispatch has_device_addr\\(b\\) is_device_ptr\\(a\\)" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(b" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "f \\(a, D\\.\[0-9\]+\\);" 1 "gimple" } }
diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c
new file mode 100644
index 0000000..ee4feff
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c
@@ -0,0 +1,5 @@
+void f(int *);
+#pragma omp declare variant(f) adjust_args(need_device_addr: x)
+/* { dg-error "expected 'nothing' or 'need_device_ptr'" "" { target *-*-* } .-1 }  */
+/* { dg-note "'need_device_addr' is not valid for C" "" { target *-*-* } .-2 }  */
+void g(int *x);