This simple patch improves the diagnostic and prepares for some future work.
Note: has_device_addr for C is permitted but pointless, for C++ it
requires some follow-up work to be useful.
For C, a hint that 'need_device_addr' it not valid has been added and
for 'has_device_addr' a middle-end warning has been added if the user
writes inconsistent code.
Comments, remarks, suggestions before I commit it?
BTW: For C++, I think we mishandle *reference to pointer type*, both by
permitting invalid code and producing and producing wrong code. But I
still have to fully understand what GCC currently does, what the spec
says it should do, and whether we want to permit a bit more for
legacy-support reasons (for this to check: GCC's testcases + older
OpenMP specifications). The topic relates to need_device_{addr,ptr},
use_device_{addr,ptr}, is_device_ptr and has_device_addr in C++, only.
Tobias
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
affect. 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' and dispatch haveing an 'has_device_addr' (or
need_device_addr/is_device_ptr) as according to the spec, those are completely
separate. Namely, '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.
gcc/c/c-parser.cc | 4 ++
gcc/cp/parser.cc | 1 +
gcc/gimplify.cc | 75 ++++++++++++++++------
gcc/testsuite/c-c++-common/gomp/adjust-args-3.c | 85 +++++++++++++++++++++++++
gcc/testsuite/gcc.dg/gomp/adjust-args-2.c | 5 ++
5 files changed, 152 insertions(+), 18 deletions(-)
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 4ec0ee85ac4..d0235809fb3 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 8a3472a4b34..15a5253b50d 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 8495c45eddb..dd0d992a958 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 00000000000..f62272cfb01
--- /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 00000000000..ee4feffb2aa
--- /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);
