The attached patch adds some more missing pieces of support for the
OpenMP "declare variant" directive -- handling for the "need_device_ptr"
and "need_device_addr" modifiers to the "adjust_args" clause. It
depends on the patch waffl3x posted last week here:
https://gcc.gnu.org/pipermail/gcc-patches/2025-April/681806.html
I've already pushed a very lightly modified backport of the current
patch to the OG14 branch, but this is the mainline version and we would
naturally like to get all of this stuff approved for mainline. :-)
To give credit where it is due, this is mostly Tobias's work; he wrote
the testcases and some patch fragments, and I just made it all work.
-Sandra
From 0cd1f6905336404f5af3c2ed4a7577e657500260 Mon Sep 17 00:00:00 2001
From: Sandra Loosemore <sloosem...@baylibre.com>
Date: Sat, 26 Apr 2025 02:22:39 +0000
Subject: [PATCH] OpenMP: need_device_ptr and need_device_addr support for
adjust_args
This patch adds support for the "need_device_addr" modifier to the
"adjust args" clause for the "declare variant" directive, and
extends/re-works the support for "need_device_ptr" as well.
This patch builds on waffl3x's recently posted patch, "OpenMP: C/C++
adjust-args numeric ranges", here.
https://gcc.gnu.org/pipermail/gcc-patches/2025-April/681806.html
In C++, "need_device_addr" supports mapping reference arguments to
device pointers. In Fortran, it similarly supports arguments passed
by reference, the default for the language, in contrast to
"need_device_ptr" which is used to map arguments of c_ptr type. The
C++ support is straightforward, but Fortran has some additional
wrinkles involving arrays passed by descriptor (a new descriptor must
be constructed with a pointer to the array data which is the only part
mapped to the device), plus special cases for passing optional
arguments and a whole array instead of a reference to its first element.
gcc/cp/ChangeLog
* parser.cc (cp_finish_omp_declare_variant): Adjust error messages.
gcc/fortran/ChangeLog
* trans-openmp.cc (gfc_trans_omp_declare_variant): Disallow
polymorphic and optional arguments with need_device_addr for now, but
don't reject need_device_addr entirely.
gcc/ChangeLog
* gimplify.cc (modify_call_for_omp_dispatch): Rework logic for
need_device_ptr and need_device_addr adjustments.
gcc/testsuite/Changelog
* c-c++-common/gomp/adjust-args-10.c: Ignore the new sorry since the
lack of proper diagnostic is already xfail'ed.
* g++.dg/gomp/adjust-args-1.C: Adjust output patterns.
* g++.dg/gomp/adjust-args-17.C: New.
* gcc.dg/gomp/adjust-args-3.c: New.
* gfortran.dg/gomp/adjust-args-14.f90: Don't expect this to fail now.
libgomp/ChangeLog
* libgomp.texi: Mark need_device_addr as supported.
* testsuite/libgomp.c++/need-device-ptr.C: New.
* testsuite/libgomp.c-c++-common/dispatch-3.c: New.
* testsuite/libgomp.fortran/adjust-args-array-descriptor.f90: New.
* testsuite/libgomp.fortran/need-device-ptr.f90: New.
Co-Authored-By: Tobias Burnus <tbur...@baylibre.com>
---
gcc/cp/parser.cc | 7 +-
gcc/fortran/trans-openmp.cc | 44 +++--
gcc/gimplify.cc | 88 +++++++--
.../c-c++-common/gomp/adjust-args-10.c | 2 +
gcc/testsuite/g++.dg/gomp/adjust-args-1.C | 6 +-
gcc/testsuite/g++.dg/gomp/adjust-args-17.C | 44 +++++
gcc/testsuite/gcc.dg/gomp/adjust-args-3.c | 47 +++++
.../gfortran.dg/gomp/adjust-args-14.f90 | 2 +-
libgomp/libgomp.texi | 2 +-
.../testsuite/libgomp.c++/need-device-ptr.C | 175 ++++++++++++++++++
.../libgomp.c-c++-common/dispatch-3.c | 35 ++++
.../adjust-args-array-descriptor.f90 | 89 +++++++++
.../libgomp.fortran/need-device-ptr.f90 | 132 +++++++++++++
13 files changed, 633 insertions(+), 40 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/adjust-args-17.C
create mode 100644 gcc/testsuite/gcc.dg/gomp/adjust-args-3.c
create mode 100644 libgomp/testsuite/libgomp.c++/need-device-ptr.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/dispatch-3.c
create mode 100644 libgomp/testsuite/libgomp.fortran/adjust-args-array-descriptor.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/need-device-ptr.f90
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 92f30d63b7e..8709b0c6181 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -50967,7 +50967,8 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
else
{
error_at (adjust_op_tok->location,
- "expected %<nothing%> or %<need_device_ptr%>");
+ "expected %<nothing%>, %<need_device_ptr%> or "
+ "%<need_device_addr%>");
/* We should be trying to recover here instead of immediately
failing, skipping to close paren and continuing. */
goto fail;
@@ -50978,8 +50979,8 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
/* We should be trying to recover here instead of immediately
failing, skipping to close paren and continuing. */
error_at (adjust_op_tok->location,
- "expected %<nothing%> or %<need_device_ptr%> followed "
- "by %<:%>");
+ "expected %<nothing%>, %<need_device_ptr%> or "
+ "%<need_device_addr%> followed by %<:%>");
goto fail;
}
/* cp_parser_omp_var_list_no_open used to handle this, we don't use
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 5fa067d9b5d..43d6b12221e 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -10025,6 +10025,34 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns, gfc_namespace *parent_ns)
&arg->sym->declared_at, &loc);
continue;
}
+ if (arg_list->u.adj_args.need_addr
+ && arg->sym->ts.type == BT_CLASS)
+ {
+ // In OpenMP 6.1, mapping polymorphic variables
+ // is undefined behavior. 'sorry' would be an
+ // alternative or some other wording.
+ gfc_error ("Argument %qs at %L to list item in "
+ "%<need_device_addr%> at %L must not "
+ "be polymorphic",
+ arg->sym->name,
+ &arg->sym->declared_at, &loc);
+ continue;
+ }
+ if (arg_list->u.adj_args.need_addr
+ && arg->sym->attr.optional)
+ {
+ // OPTIONAL has the issue that we need to handle
+ // absent arguments on the caller side, which
+ // adds extra complications.
+ gfc_error ("Sorry, argument %qs at %L to list "
+ "item in %<need_device_addr%> at %L "
+ "with OPTIONAL argument is "
+ "not yet supported",
+ arg->sym->name,
+ &arg->sym->declared_at, &loc);
+ continue;
+ }
+
if (adjust_args_list.contains (arg->sym))
{
gfc_error ("%qs at %L is specified more than "
@@ -10033,22 +10061,6 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns, gfc_namespace *parent_ns)
}
adjust_args_list.safe_push (arg->sym);
- if (arg_list->u.adj_args.need_addr)
- {
- /* TODO: Has to to support OPTIONAL and array
- descriptors; should check for CLASS, coarrays?
- Reject "abc" and 123 as actual arguments (in
- gimplify.cc or in the FE? Reject noncontiguous
- actuals? Cf. also PR C++/118859.
- Also check array-valued type(c_ptr). */
- static bool warned = false;
- if (!warned)
- sorry_at (gfc_get_location (&loc),
- "%<need_device_addr%> not yet "
- "supported");
- warned = true;
- continue;
- }
if (arg_list->u.adj_args.need_ptr
|| arg_list->u.adj_args.need_addr)
{
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 7c0854021de..db8a00db9b9 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -4413,25 +4413,81 @@ modify_call_for_omp_dispatch (tree expr, tree dispatch_clauses,
// device_num)
// but arg has to be the actual pointer, not a
// reference or a conversion expression.
- tree actual_ptr = TREE_CODE (arg) == ADDR_EXPR ? TREE_OPERAND (arg, 0)
- : arg;
- if (TREE_CODE (actual_ptr) == NOP_EXPR
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (actual_ptr, 0)))
- == REFERENCE_TYPE))
- {
- actual_ptr = TREE_OPERAND (actual_ptr, 0);
- actual_ptr
- = build1 (INDIRECT_REF, TREE_TYPE (actual_ptr), actual_ptr);
- }
tree fn = builtin_decl_explicit (BUILT_IN_OMP_GET_MAPPED_PTR);
- tree mapped_arg
- = build_call_expr_loc (loc, fn, 2, actual_ptr, dispatch_device_num);
+ tree mapped_arg = NULL_TREE;
+ bool reference_to_ptr_p = false;
- if (TREE_CODE (arg) == ADDR_EXPR
- || (TREE_CODE (TREE_TYPE (actual_ptr)) == REFERENCE_TYPE))
+ tree argtype = TREE_TYPE (arg);
+ if (!POINTER_TYPE_P (argtype))
+ {
+ sorry_at (EXPR_LOCATION (arg),
+ "Invalid non-pointer/reference argument "
+ "not diagnosed properly earlier");
+ return arg;
+ }
+
+ /* Fortran C_PTR passed by reference? Also handle the weird case
+ where an array of C_PTR is passed instead of its first element. */
+ if (need_device_ptr
+ && lang_GNU_Fortran ()
+ && (POINTER_TYPE_P (TREE_TYPE (argtype))
+ || (TREE_CODE (TREE_TYPE (argtype)) == ARRAY_TYPE
+ && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (argtype))))))
+ reference_to_ptr_p = true;
+
+ /* C++ pointer passed by reference? */
+ else if (need_device_ptr
+ && TREE_CODE (argtype) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (argtype)) == POINTER_TYPE)
+ reference_to_ptr_p = true;
+
+ /* If reference_to_ptr_p is true, we need to dereference arg to
+ get the actual pointer. */
+ tree actual_ptr = (reference_to_ptr_p
+ ? build_fold_indirect_ref (arg) : arg);
+ tree actual_ptr_type = TREE_TYPE (actual_ptr);
+ STRIP_NOPS (actual_ptr);
+
+ if (lang_hooks.decls.omp_array_data (actual_ptr, true))
+ {
+ /* This is a Fortran array with a descriptor. The actual_ptr that
+ lives on the target is the array data, not the descriptor. */
+ tree array_data
+ = lang_hooks.decls.omp_array_data (actual_ptr, false);
+ tree mapped_array_data =
+ build_call_expr_loc (loc, fn, 2, array_data, dispatch_device_num);
+
+ gcc_assert (TREE_CODE (array_data) == COMPONENT_REF);
+
+ /* We need to create a new array descriptor newd that points at the
+ mapped actual_ptr instead of the original one. Start by
+ creating the new descriptor and copy-initializing it from the
+ existing one. */
+ tree oldd = TREE_OPERAND (array_data, 0);
+ tree newd = create_tmp_var (TREE_TYPE (oldd), get_name (oldd));
+ tree t2 = build2 (MODIFY_EXPR, void_type_node, newd, oldd);
+ if (init_code)
+ init_code = build2 (COMPOUND_EXPR, void_type_node, init_code, t2);
+ else
+ init_code = t2;
+
+ /* Now stash the mapped array pointer in the new descriptor newd. */
+ tree lhs = build3 (COMPONENT_REF, TREE_TYPE (array_data), newd,
+ TREE_OPERAND (array_data, 1),
+ TREE_OPERAND (array_data, 2));
+ t2 = build2 (MODIFY_EXPR, void_type_node, lhs, mapped_array_data);
+ init_code = build2 (COMPOUND_EXPR, void_type_node, init_code, t2);
+ mapped_arg = build_fold_addr_expr (newd);
+ }
+ else
+ mapped_arg
+ = build_call_expr_loc (loc, fn, 2, actual_ptr, dispatch_device_num);
+
+ /* Cast mapped_arg back to its original type, and if we need a
+ reference, build one. */
+ mapped_arg = build1 (NOP_EXPR, actual_ptr_type, mapped_arg);
+ if (reference_to_ptr_p)
mapped_arg = build_fold_addr_expr (mapped_arg);
- else if (TREE_CODE (arg) == NOP_EXPR)
- mapped_arg = build1 (NOP_EXPR, TREE_TYPE (arg), mapped_arg);
return mapped_arg;
};
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-10.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-10.c
index 5cda21e07ee..6730dfeba2d 100644
--- a/gcc/testsuite/c-c++-common/gomp/adjust-args-10.c
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-10.c
@@ -11,3 +11,5 @@ void f0(int *p0, int *p1, int *p2, int *p3, int *p4)
#pragma omp dispatch
b0(p0, p1, p2, p3, p4, 42); /* { dg-error "variadic argument 5 specified in an 'append_args' clause with the 'need_device_ptr' modifier must be of pointer type" "" { xfail *-*-* } } */
}
+
+/* { dg-prune-output "sorry, unimplemented: Invalid non-pointer/reference argument not diagnosed properly earlier" } */
diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-1.C b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C
index 3aee78e3bb3..d0e0bce7444 100644
--- a/gcc/testsuite/g++.dg/gomp/adjust-args-1.C
+++ b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C
@@ -13,13 +13,13 @@ int f2a (void *a);
int f2b (void *a);
#pragma omp declare variant (f0) match (construct={dispatch},device={arch(gcn)}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */
int f2c (void *a);
-#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr'" } */
int f3 (int a);
#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause requires a 'match' clause" } */
int f4 (void *a);
-#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" } */
int f5 (int a);
-#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" } */
int f6 (int a);
#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected primary-expression before '\\)' token" } */
int f7 (int a);
diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-17.C b/gcc/testsuite/g++.dg/gomp/adjust-args-17.C
new file mode 100644
index 00000000000..62ddab0b74b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/adjust-args-17.C
@@ -0,0 +1,44 @@
+void f(int*,int &,int*);
+void f0(int*,int &,int*);
+void f1(int*,int &,int*);
+void f2(int*,int &,int*);
+void f3(int*,int &,int*);
+void f4(int*,int &,int*);
+void f5(int*,int &,int*);
+void f6(int*,int &,int*);
+void f7(int*,int &,int*);
+void f8(int*,int &,int*);
+void f9(int*,int &,int*);
+void fa(int*,int &,int*);
+void f10(int*,int &,int*);
+void f11(int*,int &,int*);
+void f12(int*,int &,int*);
+void f13(int*,int &,int*);
+void f14(int*,int &,int*);
+void f15(int*,int &,int*);
+void f16(int*,int &,int*);
+
+#pragma omp declare variant(f) match(construct={dispatch}) adjust_args(x : y) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr'" }
+#pragma omp declare variant(f0) match(construct={dispatch}) adjust_args(x) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f1) match(construct={dispatch}) adjust_args(x,) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f2) match(construct={dispatch}) adjust_args(foo x) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f3) match(construct={dispatch}) adjust_args(nothing) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f4) match(construct={dispatch}) adjust_args(need_device_ptr) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f5) match(construct={dispatch}) adjust_args(nothing x) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f6) match(construct={dispatch}) adjust_args(need_device_ptr x) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f7) match(construct={dispatch}) adjust_args(need_device_addr x) // { dg-error "expected 'nothing', 'need_device_ptr' or 'need_device_addr' followed by ':'" }
+#pragma omp declare variant(f8) match(construct={dispatch}) adjust_args(nothing :) // { dg-error "expected primary-expression before '\\)' token" }
+#pragma omp declare variant(f9) match(construct={dispatch}) adjust_args(need_device_ptr :) // { dg-error "expected primary-expression before '\\)' token" }
+#pragma omp declare variant(fa) match(construct={dispatch}) adjust_args(need_device_addr :) // { dg-error "expected primary-expression before '\\)' token" }
+#pragma omp declare variant(f10) match(construct={dispatch}) adjust_args(need_device_addr : omp_num_args-1) // { dg-error "expected ':' before '\\)' token" }
+// { dg-note "93: an expression is only allowed in a numeric range" "" { target *-*-* } .-1 }
+
+// Valid:
+#pragma omp declare variant(f11) match(construct={dispatch}) adjust_args(nothing : z, 1:2)
+#pragma omp declare variant(f12) match(construct={dispatch}) adjust_args(need_device_ptr : x)
+#pragma omp declare variant(f13) match(construct={dispatch}) adjust_args(need_device_addr : y)
+#pragma omp declare variant(f14) match(construct={dispatch}) adjust_args(nothing : :)
+#pragma omp declare variant(f15) match(construct={dispatch}) adjust_args(need_device_ptr : 3:3)
+#pragma omp declare variant(f16) match(construct={dispatch}) adjust_args(need_device_addr : 2:2)
+
+void g(int*x, int &y, int *z);
diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-3.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-3.c
new file mode 100644
index 00000000000..a9e7fabab2b
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-3.c
@@ -0,0 +1,47 @@
+void f(int*,int *,int*);
+void f0(int*,int *,int*);
+void f1(int*,int *,int*);
+void f2(int*,int *,int*);
+void f3(int*,int *,int*);
+void f4(int*,int *,int*);
+void f5(int*,int *,int*);
+void f6(int*,int *,int*);
+void f7(int*,int *,int*);
+void f8(int*,int *,int*);
+void f9(int*,int *,int*);
+void fa(int*,int *,int*);
+void f10(int*,int *,int*);
+void f11(int*,int *,int*);
+void f12(int*,int *,int*);
+void f13(int*,int *,int*);
+void f14(int*,int *,int*);
+void f15(int*,int *,int*);
+void f16(int*,int *,int*);
+
+#pragma omp declare variant(f) match(construct={dispatch}) adjust_args(x : y) // { dg-error "expected 'nothing' or 'need_device_ptr'" }
+#pragma omp declare variant(f0) match(construct={dispatch}) adjust_args(x) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f1) match(construct={dispatch}) adjust_args(x,) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f2) match(construct={dispatch}) adjust_args(foo x) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f3) match(construct={dispatch}) adjust_args(nothing) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f4) match(construct={dispatch}) adjust_args(need_device_ptr) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f5) match(construct={dispatch}) adjust_args(nothing x) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f6) match(construct={dispatch}) adjust_args(need_device_ptr x) // { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" }
+#pragma omp declare variant(f7) match(construct={dispatch}) adjust_args(need_device_addr x) // { dg-error "expected 'nothing' or 'need_device_ptr'" }
+#pragma omp declare variant(f8) match(construct={dispatch}) adjust_args(nothing :) // { dg-error "expected expression before '\\)' token" }
+#pragma omp declare variant(f9) match(construct={dispatch}) adjust_args(need_device_ptr :) // { dg-error "expected expression before '\\)' token" }
+#pragma omp declare variant(fa) match(construct={dispatch}) adjust_args(need_device_addr :) // { dg-error "expected 'nothing' or 'need_device_ptr'" }
+// { dg-note "73: 'need_device_addr' is not valid for C" "" { target *-*-* } .-1 }
+#pragma omp declare variant(f10) match(construct={dispatch}) adjust_args(need_device_ptr : omp_num_args-1) // { dg-error "expected ':' before '\\)' token" }
+// { dg-note "92: an expression is only allowed in a numeric range" "" { target *-*-* } .-1 }
+
+// Valid:
+#pragma omp declare variant(f11) match(construct={dispatch}) adjust_args(nothing : z, 1:2)
+#pragma omp declare variant(f12) match(construct={dispatch}) adjust_args(need_device_ptr : x)
+#pragma omp declare variant(f13) match(construct={dispatch}) adjust_args(need_device_addr : y) // { dg-error "expected 'nothing' or 'need_device_ptr'" }
+// { dg-note "74: 'need_device_addr' is not valid for C" "" { target *-*-* } .-1 }
+#pragma omp declare variant(f14) match(construct={dispatch}) adjust_args(nothing : :)
+#pragma omp declare variant(f15) match(construct={dispatch}) adjust_args(need_device_ptr : 3:3)
+#pragma omp declare variant(f16) match(construct={dispatch}) adjust_args(need_device_addr : 2:2)// { dg-error "expected 'nothing' or 'need_device_ptr'" }
+// { dg-note "74: 'need_device_addr' is not valid for C" "" { target *-*-* } .-1 }
+
+void g(int*x, int *y, int *z);
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-14.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-14.f90
index e644fd7060e..95b039ef497 100644
--- a/gcc/testsuite/gfortran.dg/gomp/adjust-args-14.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-14.f90
@@ -14,7 +14,7 @@ contains
! { dg-error "19: Argument 'y' at .1. to list item in 'need_device_addr' at .2. must not have the VALUE attribute" "" { target *-*-* } 8 }
! { dg-error "62: Argument 'y' at .1. to list item in 'need_device_addr' at .2. must not have the VALUE attribute" "" { target *-*-* } 9 }
-! { dg-message "sorry, unimplemented: 'need_device_addr' not yet supported" "" { target *-*-* } 9 }
+
! { dg-error "Argument 'z' at .1. to list item in 'need_device_ptr' at .2. must be a scalar of TYPE\\(C_PTR\\)" "" { target *-*-* } 8 }
! { dg-error "Argument 'z' at .1. to list item in 'need_device_ptr' at .2. must be a scalar of TYPE\\(C_PTR\\)" "" { target *-*-* } 10 }
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index d57afac859c..2893896f587 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -543,7 +543,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
@item New @code{partitioner} value to @code{partition} allocator trait
@tab N @tab
@item Semicolon-separated list to @code{uses_allocators} @tab N @tab
-@item New @code{need_device_addr} modifier to @code{adjust_args} clause @tab N @tab
+@item New @code{need_device_addr} modifier to @code{adjust_args} clause @tab Y @tab
@item @code{interop} clause to @code{dispatch} @tab Y @tab
@item Scope requirement changes for @code{declare_target} @tab N @tab
@item @code{message} and @code{severity} clauses to @code{parallel} directive
diff --git a/libgomp/testsuite/libgomp.c++/need-device-ptr.C b/libgomp/testsuite/libgomp.c++/need-device-ptr.C
new file mode 100644
index 00000000000..d7babffae96
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/need-device-ptr.C
@@ -0,0 +1,175 @@
+// Test the need_device_ptr and need_device_addr modifiers to the adjust_args clause
+
+#include <omp.h>
+
+void fptr_var (int *x1, int *x2, int *x3, int **x3a, int *x4, int *x5, int *x6, int **x6a)
+{
+ #pragma omp target is_device_ptr (x1)
+ { if (*x1 != 1) __builtin_abort (); *x1 *= -1; }
+
+ #pragma omp target is_device_ptr (x2)
+ { if (*x2 != 2) __builtin_abort (); *x2 *= -1; }
+
+ #pragma omp target is_device_ptr (x3)
+ { if (*x3 != 3) __builtin_abort (); *x3 *= -1; }
+
+ #pragma omp target is_device_ptr (x3a)
+ { if (**x3a != 30) __builtin_abort (); **x3a *= -1; }
+
+ #pragma omp target is_device_ptr (x4)
+ { if (*x4 != 4) __builtin_abort (); *x4 *= -1; }
+
+ #pragma omp target is_device_ptr (x5)
+ { if (*x5 != 5) __builtin_abort (); *x5 *= -1; }
+
+ #pragma omp target is_device_ptr (x6)
+ { if (*x6 != 6) __builtin_abort (); *x6 *= -1; }
+
+ #pragma omp target is_device_ptr (x6a)
+ { if (**x6a != 60) __builtin_abort (); **x6a *= -1; }
+}
+
+#pragma omp declare variant(fptr_var) match(construct={dispatch}) adjust_args (need_device_ptr : 1:8)
+void fptr (int *x1, int *x2, int *x3, int **x3a, int *x4, int *x5, int *x6, int **x6a);
+
+void faddr_var (int &x1, int &x2, int &x3, int *&x3a, int &x4, int &x5, int &x6, int *&x6a)
+{
+ #pragma omp target has_device_addr (x1)
+ { if (x1 != 1) __builtin_abort (); x1 *= -1; }
+
+ #pragma omp target has_device_addr (x2)
+ { if (x2 != 2) __builtin_abort (); x2 *= -1; }
+
+ #pragma omp target has_device_addr (x3)
+ { if (x3 != 3) __builtin_abort (); x3 *= -1; }
+
+ #pragma omp target has_device_addr (x3a)
+ { if (*x3a != 30) __builtin_abort (); *x3a *= -1; }
+
+ #pragma omp target has_device_addr (x4)
+ { if (x4 != 4) __builtin_abort (); x4 *= -1; }
+
+ #pragma omp target has_device_addr (x5)
+ { if (x5 != 5) __builtin_abort (); x5 *= -1; }
+
+ #pragma omp target has_device_addr (x6)
+ { if (x6 != 6) __builtin_abort (); x6 *= -1; }
+
+ #pragma omp target has_device_addr (x6a)
+ { if (*x6a != 60) __builtin_abort (); *x6a *= -1; }
+}
+
+#pragma omp declare variant(faddr_var) match(construct={dispatch}) adjust_args (need_device_addr : 1:8)
+void faddr (int &x1, int &x2, int &x3, int *&, int &x4, int &x5, int &x6, int *&);
+
+void caller_ptr(int x, int &y, int *z, int *zptr)
+{
+ int a = 4;
+ int bval = 5;
+ int &b = bval;
+ int *c = (int*) __builtin_malloc (sizeof (int));
+ int *cptr;
+ *c = 6;
+
+ zptr = (int *) omp_target_alloc (sizeof (int), omp_get_default_device ());
+ cptr = (int *) omp_target_alloc (sizeof (int), omp_get_default_device ());
+
+ #pragma omp target is_device_ptr(cptr, zptr)
+ {
+ *zptr = 30;
+ *cptr = 60;
+ }
+
+ #pragma omp target enter data map(x, a, b, c[:1], cptr, zptr)
+
+ #pragma omp dispatch
+ fptr (&x, &y, z, &zptr, &a, &b, c, &cptr);
+
+ #pragma omp target exit data map(x, a, b, c[:1], cptr, zptr)
+ #pragma omp target update from(y, z[:1])
+
+ if (x != -1) __builtin_abort ();
+ if (y != -2) __builtin_abort ();
+ if (*z != -3) __builtin_abort ();
+
+ if (a != -4) __builtin_abort ();
+ if (b != -5) __builtin_abort ();
+ if (*c != -6) __builtin_abort ();
+
+ #pragma omp target is_device_ptr(cptr, zptr)
+ {
+ if (*zptr != -30) __builtin_abort ();
+ if (*cptr != -60) __builtin_abort ();
+ }
+
+ __builtin_free (c);
+ omp_target_free (cptr, omp_get_default_device ());
+ omp_target_free (zptr, omp_get_default_device ());
+}
+
+void caller_addr(int x, int &y, int *z, int *zptr)
+{
+ int a = 4;
+ int bval = 5;
+ int &b = bval;
+ int *c = (int*) __builtin_malloc (sizeof (int));
+ int *cptr;
+ *c = 6;
+
+ zptr = (int *) omp_target_alloc (sizeof (int), omp_get_default_device ());
+ cptr = (int *) omp_target_alloc (sizeof (int), omp_get_default_device ());
+
+ #pragma omp target is_device_ptr(cptr, zptr)
+ {
+ *zptr = 30;
+ *cptr = 60;
+ }
+
+ #pragma omp target enter data map(x, a, b, c[:1], cptr, zptr)
+
+ #pragma omp dispatch
+ faddr (x, y, *z, zptr, a, b, *c, cptr);
+
+ #pragma omp target exit data map(x, a, b, c[:1], cptr, zptr)
+ #pragma omp target update from(y, z[:1])
+
+ if (x != -1) __builtin_abort ();
+ if (y != -2) __builtin_abort ();
+ if (*z != -3) __builtin_abort ();
+
+ if (a != -4) __builtin_abort ();
+ if (b != -5) __builtin_abort ();
+ if (*c != -6) __builtin_abort ();
+
+ #pragma omp target is_device_ptr(cptr, zptr)
+ {
+ if (*zptr != -30) __builtin_abort ();
+ if (*cptr != -60) __builtin_abort ();
+ }
+
+
+ __builtin_free (c);
+}
+
+int
+main ()
+{
+ int x = 1;
+ int yval = 2;
+ int &y = yval;
+ int *z = (int *) __builtin_malloc (sizeof (int));
+ int *zptr;
+ *z = 3;
+
+ #pragma omp target data map(y, z[:1])
+ caller_ptr (x, y, z, zptr);
+
+ x = 1;
+ y = 2;
+ *z = 3;
+
+ #pragma omp target data map(y, z[:1], zptr)
+ caller_addr (x, y, z, zptr);
+
+ __builtin_free (z);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/dispatch-3.c b/libgomp/testsuite/libgomp.c-c++-common/dispatch-3.c
new file mode 100644
index 00000000000..2c41e3cd470
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/dispatch-3.c
@@ -0,0 +1,35 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+/* PR c++/118859 */
+
+void f_var(int *y) {
+ #pragma omp target is_device_ptr(y)
+ {
+ if (*y != 5)
+ __builtin_abort ();
+ *y += 10;
+ }
+}
+#pragma omp declare variant(f_var) match(construct={dispatch}) adjust_args(need_device_ptr : 1)
+void f(int *);
+
+static void test()
+{
+ int x = 5;
+ #pragma omp target enter data map(x)
+
+ #pragma omp dispatch
+ f(&x);
+
+ #pragma omp target exit data map(x)
+ if (x != 15)
+ __builtin_abort ();
+}
+
+int main()
+{
+ test();
+}
+
+// { dg-final { scan-tree-dump "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(&x, D\\.\[0-9\]+\\);" "gimple" } }
+// { dg-final { scan-tree-dump "f_var \\(D\\.\[0-9\]+\\);" "gimple" } }
diff --git a/libgomp/testsuite/libgomp.fortran/adjust-args-array-descriptor.f90 b/libgomp/testsuite/libgomp.fortran/adjust-args-array-descriptor.f90
new file mode 100644
index 00000000000..dd9b57b8387
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/adjust-args-array-descriptor.f90
@@ -0,0 +1,89 @@
+! Test array descriptor handling with the need_device_addr modifier to adjust_args
+
+module m
+ use iso_c_binding
+ implicit none (type, external)
+
+ integer :: case = 0
+contains
+ subroutine var_array_alloc(x)
+ integer, allocatable :: x(:)
+ !$omp target has_device_addr(x)
+ block
+ if (size(x) /= 3) stop 1
+ if (any (x /= [1,2,3])) stop 2
+ x = x * (-1)
+ end block
+ end
+
+ subroutine base_array_alloc(x)
+ !$omp declare variant(var_array_alloc) match(construct={dispatch}) adjust_args(need_device_addr : x)
+ integer, allocatable :: x(:)
+ error stop
+ end
+
+ subroutine var_array_nonalloc(x)
+ integer :: x(:)
+ !$omp target has_device_addr(x)
+ block
+ if (size(x) /= 4) stop 3
+ if (any (x /= [11,22,33,44])) stop 4
+ x = x * (-1)
+ end block
+ end
+
+ subroutine base_array_nonalloc(x)
+ !$omp declare variant(var_array_nonalloc) match(construct={dispatch}) adjust_args(need_device_addr : x)
+ integer :: x(:)
+ error stop
+ end
+
+ subroutine test_array_alloc(y)
+ integer, allocatable :: y(:)
+ !$omp target enter data map(y)
+
+
+ ! Direct call (for testing; value check fails if both are enabled
+ ! !$omp target data use_device_addr(y)
+ ! call var_array_alloc (y)
+ ! !$omp end target data
+
+ !$omp dispatch
+ call base_array_alloc (y)
+
+ !$omp target exit data map(y)
+
+ if (size(y) /= 3) stop 3
+ if (any (y /= [-1,-2,-3])) stop 1
+ end
+
+ subroutine test_array_nonalloc()
+ integer :: y(4)
+ y = [11,22,33,44]
+
+ !$omp target enter data map(y)
+
+ ! Direct call (for testing; value check fails if both are enabled
+ !!$omp target data use_device_addr(y)
+ ! call var_array_nonalloc (y)
+ !!$omp end target data
+
+ !$omp dispatch
+ call base_array_nonalloc (y)
+
+ !$omp target exit data map(y)
+
+ if (size(y) /= 4) stop 3
+ if (any (y /= [-11,-22,-33,-44])) stop 1
+ end
+end module
+
+use m
+implicit none
+integer, allocatable :: z(:)
+
+z = [1,2,3]
+call test_array_alloc(z)
+call test_array_nonalloc()
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/need-device-ptr.f90 b/libgomp/testsuite/libgomp.fortran/need-device-ptr.f90
new file mode 100644
index 00000000000..c75688c3486
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/need-device-ptr.f90
@@ -0,0 +1,132 @@
+! Comprehensive non-array testcase for need_device_ptr / need_device_addr
+
+module m
+ use iso_c_binding
+ implicit none (type, external)
+
+ integer :: case = 0
+contains
+ subroutine var_ptr_f(n, x, y, z)
+ integer, value :: n
+ type(c_ptr) :: x
+ type(c_ptr), value :: y
+ type(c_ptr), optional :: z
+ !$omp target is_device_ptr(x,y,z)
+ block
+ integer, pointer :: ix, iy, iz
+ call c_f_pointer(x, ix)
+ call c_f_pointer(y, iy)
+ call c_f_pointer(z, iz)
+ if (ix /= 52) stop n*10 + 1
+ if (iy /= 85) stop n*10 + 2
+ if (iz /= 52) stop n*10 + 5
+ end block
+ end
+ subroutine base_ptr_f(n, x, y, z)
+ !$omp declare variant(var_ptr_f) match(construct={dispatch}) adjust_args(need_device_ptr : x, y, z)
+ integer, value :: n
+ type(c_ptr) :: x
+ type(c_ptr), value :: y
+ type(c_ptr), optional :: z
+ error stop n
+ end
+
+ subroutine var_caddr_f(x, y)
+ type(c_ptr) :: x
+ type(c_ptr), optional :: y
+ !$omp target has_device_addr(x, y)
+ block
+ integer, pointer :: ix, iy
+ call c_f_pointer(x, ix)
+ call c_f_pointer(x, iy)
+ if (ix /= 52) stop 3
+ if (iy /= 85) stop 6
+ end block
+ end
+! FIXME: optional args give a "sorry".
+! subroutine base_caddr_f(x, y)
+! !$omp declare variant(var_caddr_f) match(construct={dispatch}) adjust_args(need_device_addr : x, y)
+! type(c_ptr) :: x
+! type(c_ptr), optional :: y
+! error stop
+! end
+
+ subroutine var_iaddr_f(x,y)
+ integer :: x
+ integer, optional :: y
+ !$omp target has_device_addr(x, y)
+ block
+ if (x /= 52) stop 4
+ if (y /= 85) stop 4
+ end block
+ end
+
+! FIXME: optional args give a "sorry".
+! subroutine base_iaddr_f(x,y)
+! !$omp declare variant(var_iaddr_f) match(construct={dispatch}) adjust_args(need_device_addr : x, y)
+! integer :: x
+! integer, optional :: y
+! error stop
+! end
+
+ subroutine test_f(carg1, carg2, carg1v, carg2v, iarg1, iarg2)
+ type(c_ptr) :: carg1, carg2
+ type(c_ptr), value :: carg1v, carg2v
+ integer, target :: iarg1, iarg2
+ type(c_ptr) :: cptr1, cptr2
+ integer, target :: ivar1, ivar2
+
+
+ ivar1 = 52
+ ivar2 = 85
+
+ !$omp target enter data map(to: ivar1, ivar2)
+
+ cptr1 = c_loc(ivar1)
+ cptr2 = c_loc(ivar2)
+
+ !$omp dispatch
+ call base_ptr_f (1, carg1, carg2, carg1)
+ !$omp dispatch
+ call base_ptr_f (2, carg1v, carg2v, carg1v)
+ !$omp dispatch
+ call base_ptr_f (3, cptr1, cptr2, cptr1)
+ !$omp dispatch
+ call base_ptr_f (4, c_loc(iarg1), c_loc(iarg2), c_loc(iarg1))
+ !$omp dispatch
+ call base_ptr_f (6, c_loc(ivar1), c_loc(ivar2), c_loc(ivar1))
+
+! FIXME: optional argument functions not supported yet.
+! !$omp dispatch
+! call base_caddr_f (carg1, carg2)
+! !$omp dispatch
+! call base_caddr_f (carg1v, carg2v)
+! !$omp dispatch
+! call base_caddr_f (cptr1, cptr2)
+! !$omp dispatch
+! call base_caddr_f (c_loc(iarg1), c_loc(iarg2))
+! !$omp dispatch
+! call base_caddr_f (c_loc(ivar1), c_loc(ivar2))
+! !$omp dispatch
+! call base_iaddr_f (iarg1, iarg2)
+! !$omp dispatch
+! call base_iaddr_f (ivar1, iarg2)
+
+ !$omp target exit data map(release: ivar1, ivar2)
+ end
+end module m
+
+use m
+implicit none
+integer, target :: mx, my
+type(c_ptr) :: cptr1, cptr2
+mx = 52
+my = 85
+
+cptr1 = c_loc(mx)
+cptr2 = c_loc(my)
+
+!$omp target data map(to: mx, my)
+ call test_f (cptr1, cptr2, cptr1, cptr2, mx, my)
+!$omp end target data
+end
--
2.34.1
