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.

1 files changed, 85 insertions, 0 deletions

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" } }