[gcc/devel/omp/gcc-16] OpenMP: need_device_ptr and need_device_addr support for adjust_args
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:f097b1c7f973afec58ec55a57543725481191b09 commit f097b1c7f973afec58ec55a57543725481191b09 Author: Sandra Loosemore <[email protected]> Date: Wed May 6 19:37:53 2026 +0000 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 <[email protected]> Diff: --- gcc/cp/parser.cc | 7 +- gcc/fortran/trans-openmp.cc | 44 ++++-- gcc/gimplify.cc | 88 +++++++++-- gcc/testsuite/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 ++++++ gcc/testsuite/gfortran.dg/gomp/adjust-args-14.f90 | 2 +- libgomp/libgomp.texi | 2 +- libgomp/testsuite/libgomp.c++/need-device-ptr.C | 175 +++++++++++++++++++++ .../testsuite/libgomp.c-c++-common/dispatch-3.c | 35 +++++ .../adjust-args-array-descriptor.f90 | 89 +++++++++++ .../testsuite/libgomp.fortran/need-device-ptr.f90 | 132 ++++++++++++++++ 13 files changed, 633 insertions(+), 40 deletions(-) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 67d64b2147d8..1dfaf85f519b 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -54835,7 +54835,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; @@ -54846,8 +54847,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 cfc45573ae68..7c1a6bb44274 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -11194,6 +11194,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 " @@ -11202,22 +11230,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 0eadd4477504..b959fd834ad1 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -4587,25 +4587,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; + + 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); - if (TREE_CODE (arg) == ADDR_EXPR - || (TREE_CODE (TREE_TYPE (actual_ptr)) == REFERENCE_TYPE)) + /* 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 5cda21e07eee..6730dfeba2de 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 105bfef5a29e..a7c8eef3d79c 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 "expected '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 000000000000..62ddab0b74bb --- /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 000000000000..a9e7fabab2b3 --- /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 e644fd7060e3..95b039ef4977 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 33f9e5cdbefe..0dc923c479d5 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -544,7 +544,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 000000000000..d7babffae966 --- /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 000000000000..2c41e3cd470b --- /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 000000000000..dd9b57b8387f --- /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 000000000000..c75688c34862 --- /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