https://gcc.gnu.org/g:b57c24e76e6e752901c964395b8f6803a6d9f1f5
commit b57c24e76e6e752901c964395b8f6803a6d9f1f5
Author: Tobias Burnus <tbur...@baylibre.com>
Date: Mon Jan 27 12:40:17 2025 +0100
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.
(cherry picked from commit 2cbb2408a830a63fbd901a4da3bfd341cec4b6ef)
Diff:
---
gcc/ChangeLog.omp | 9 +++
gcc/c/ChangeLog.omp | 9 +++
gcc/c/c-parser.cc | 4 ++
gcc/cp/ChangeLog.omp | 7 ++
gcc/cp/parser.cc | 1 +
gcc/gimplify.cc | 76 ++++++++++++++++------
gcc/testsuite/ChangeLog.omp | 8 +++
gcc/testsuite/c-c++-common/gomp/adjust-args-3.c | 85 +++++++++++++++++++++++++
gcc/testsuite/gcc.dg/gomp/adjust-args-2.c | 5 ++
9 files changed, 186 insertions(+), 18 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index e83fe6afebb4..36e1767e841e 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2025-01-27 Tobias Burnus <tbur...@baylibre.com>
+
+ Backported from master:
+ 2024-12-12 Tobias Burnus <tbur...@baylibre.com>
+
+ * 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.
+
2025-01-27 Tobias Burnus <tbur...@baylibre.com>
Backported from master:
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index a19c27f82634..cf0e23a0a582 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,12 @@
+2025-01-27 Tobias Burnus <tbur...@baylibre.com>
+
+ Backported from master:
+ 2024-12-12 Tobias Burnus <tbur...@baylibre.com>
+
+ * 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.
+
2025-01-27 Tobias Burnus <tbur...@baylibre.com>
Backported from master:
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 4c1cead30248..acb07c869db1 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -25049,6 +25049,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))
@@ -26780,6 +26781,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index feb5fbb0b09d..0ea0b84c51a6 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,10 @@
+2025-01-27 Tobias Burnus <tbur...@baylibre.com>
+
+ Backported from master:
+ 2024-12-12 Tobias Burnus <tbur...@baylibre.com>
+
+ * parser.cc (OMP_DISPATCH_CLAUSE_MASK): Add has_device_addr clause.
+
2025-01-27 Tobias Burnus <tbur...@baylibre.com>
Backported from master:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 7d65fc2fa674..042a5a460e3a 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -50191,6 +50191,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 cd9bb8d701e1..cfb0114c3414 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -4393,27 +4393,39 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p,
fallback_t fallback)
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
@@ -4424,15 +4436,43 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p,
fallback_t fallback)
|| 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 (variant_substituted_p && !is_device_ptr)
+ if (variant_substituted_p
+ && ((need_device_ptr && !is_device_ptr)
+ || (need_device_addr && !has_device_addr)))
{
if (device_num == NULL_TREE)
{
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index bd3d806ad08a..35da06819a3e 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,11 @@
+2025-01-27 Tobias Burnus <tbur...@baylibre.com>
+
+ Backported from master:
+ 2024-12-12 Tobias Burnus <tbur...@baylibre.com>
+
+ * c-c++-common/gomp/adjust-args-3.c: New test.
+ * gcc.dg/gomp/adjust-args-2.c: New test.
+
2025-01-27 Tobias Burnus <tbur...@baylibre.com>
Backported from master:
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 000000000000..f62272cfb019
--- /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 000000000000..ee4feffb2aa3
--- /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);
