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.

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 4ec0ee85ac498739cfaffcbd7904cc5ebfb4fc07..d0235809fb38ef05821c84c4d040c91de519df44 100644 (file)
--- 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 8a3472a4b34e4908f3e50d54cc040426206dfbaf..15a5253b50d1808fdd879a0dcd695cae3bd3a4e8 100644 (file)
--- 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 8495c45eddbb87bd20cbae42aff0ad7cd49439a7..dd0d992a9581d2dbdbbf225328004eb4f5b05cc4 100644 (file)
--- 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 (file)
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 (file)
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);