mirror of git://gcc.gnu.org/git/gcc.git
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-c++-common/dispatch-3.c: New. * testsuite/libgomp.c++/need-device-ptr.C: New. * testsuite/libgomp.fortran/adjust-args-array-descriptor.f90: New. * testsuite/libgomp.fortran/need-device-ptr.f90: New. Co-Authored-By: Tobias Burnus <tburnus@baylibre.com>
This commit is contained in:
Sandra Loosemore
parent
97067daf1d
commit
9a06e4d6a1
|
|
@ -51407,7 +51407,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;
|
||||
|
|
@ -51418,8 +51419,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
|
||||
|
|
|
|||
|
|
@ -11968,6 +11968,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 "
|
||||
|
|
@ -11976,22 +12004,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)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -4434,25 +4434,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;
|
||||
};
|
||||
|
||||
|
|
|
|||
|
|
@ -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" } */
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
@ -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);
|
||||
|
|
@ -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 "64: 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 }
|
||||
|
|
|
|||
|
|
@ -516,6 +516,7 @@ Technical Report (TR) 12 is the second preview for OpenMP 6.0.
|
|||
@item Extension of @code{interop} operation of @code{append_args}, allowing all
|
||||
modifiers of the @code{init} 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 @code{message} and @code{severity} clauses to @code{parallel} directive
|
||||
@tab N @tab
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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" } }
|
||||
|
|
@ -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
|
||||
|
|
@ -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
|
||||
Loading…
Reference in New Issue