public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Kwok Yeung <kcy@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] openmp: Add warning when functions containing metadirectives with 'construct={target}' called direct Date: Wed, 29 Jun 2022 14:44:09 +0000 (GMT) [thread overview] Message-ID: <20220629144409.58049384D15E@sourceware.org> (raw) https://gcc.gnu.org/g:2e950ecaeaac966f2a717d32e4c6a8ad4f2661d2 commit 2e950ecaeaac966f2a717d32e4c6a8ad4f2661d2 Author: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Fri Jan 28 13:56:33 2022 +0000 openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly void f(void) { #pragma omp metadirective \ when (construct={target}: A) \ default (B) ... } ... { #pragma omp target f(); // Target call f(); // Local call } With the OpenMP 5.0/5.1 specifications, we would expect A to be selected in the metadirective when the target call is made, but B when f is called directly outside of a target context. However, since GCC does not have separate copies of f for local and target calls, and the construct selector is static, it must be resolved one way or the other at compile-time (currently in the favour of selecting A), which may be unexpected behaviour. This patch attempts to detect the above situation, and will emit a warning if found. 2022-01-28 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/ * gimplify.cc (gimplify_omp_metadirective): Mark offloadable functions containing metadirectives with 'construct={target}' in the selector. * omp-general.cc (omp_has_target_constructor_p): New. * omp-general.h (omp_has_target_constructor_p): New prototype. * omp-low.cc (lower_omp_1): Emit warning if marked functions called outside of a target context. gcc/testsuite/ * c-c++-common/gomp/metadirective-4.c (main): Add expected warning. * gfortran.dg/gomp/metadirective-4.f90 (test): Likewise. libgomp/ * testsuite/libgomp.c-c++-common/metadirective-2.c (main): Add expected warning. * testsuite/libgomp.fortran/metadirective-2.f90 (test): Likewise. Diff: --- gcc/ChangeLog.omp | 9 +++++++++ gcc/gimplify.cc | 21 +++++++++++++++++++++ gcc/omp-general.cc | 21 +++++++++++++++++++++ gcc/omp-general.h | 1 + gcc/omp-low.cc | 18 ++++++++++++++++++ gcc/testsuite/ChangeLog.omp | 5 +++++ gcc/testsuite/c-c++-common/gomp/metadirective-4.c | 2 +- gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 | 2 +- libgomp/ChangeLog.omp | 6 ++++++ .../libgomp.c-c++-common/metadirective-2.c | 2 +- .../testsuite/libgomp.fortran/metadirective-2.f90 | 2 +- 11 files changed, 85 insertions(+), 4 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index eb3a1130f72..abcdb020357 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,12 @@ +2022-01-28 Kwok Cheung Yeung <kcy@codesourcery.com> + + * gimplify.cc (gimplify_omp_metadirective): Mark offloadable functions + containing metadirectives with 'construct={target}' in the selector. + * omp-general.cc (omp_has_target_constructor_p): New. + * omp-general.h (omp_has_target_constructor_p): New prototype. + * omp-low.cc (lower_omp_1): Emit warning if marked functions called + outside of a target context. + 2022-01-25 Kwok Cheung Yeung <kcy@codesourcery.com> * builtin-types.def (BT_FN_BOOL_INT_CONST_PTR_CONST_PTR_CONST_PTR): New diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index fca7f4d6bf2..105bd81b5b8 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -15629,6 +15629,27 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *, { auto_vec<tree> selectors; + /* Mark offloadable functions containing metadirectives that specify + a 'construct' selector with a 'target' constructor. */ + if (offloading_function_p (current_function_decl)) + { + for (tree clause = OMP_METADIRECTIVE_CLAUSES (*expr_p); + clause != NULL_TREE; clause = TREE_CHAIN (clause)) + { + tree selector = TREE_PURPOSE (clause); + + if (omp_has_target_constructor_p (selector)) + { + tree id = get_identifier ("omp metadirective construct target"); + + DECL_ATTRIBUTES (current_function_decl) + = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (current_function_decl)); + break; + } + } + } + /* Try to resolve the metadirective. */ vec<struct omp_metadirective_variant> candidates = omp_resolve_metadirective (*expr_p); diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 48953e11997..05b960ea983 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -2939,6 +2939,27 @@ omp_resolve_metadirective (gimple *gs) return omp_get_dynamic_candidates (variants); } +bool +omp_has_target_constructor_p (tree selector) +{ + if (selector == NULL_TREE) + return false; + + tree selector_set = TREE_PURPOSE (selector); + if (strcmp (IDENTIFIER_POINTER (selector_set), "construct") != 0) + return false; + + enum tree_code constructs[5]; + int nconstructs + = omp_constructor_traits_to_codes (TREE_VALUE (selector), constructs); + + for (int i = 0; i < nconstructs; i++) + if (constructs[i] == OMP_TARGET) + return true; + + return false; +} + /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK macro on gomp-constants.h. We do not check for overflow. */ diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 65c1a60ad41..6074b0214f4 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -126,6 +126,7 @@ extern tree omp_get_context_selector (tree, const char *, const char *); extern tree omp_resolve_declare_variant (tree); extern vec<struct omp_metadirective_variant> omp_resolve_metadirective (tree); extern vec<struct omp_metadirective_variant> omp_resolve_metadirective (gimple *); +extern bool omp_has_target_constructor_p (tree); extern tree oacc_launch_pack (unsigned code, tree device, unsigned op); extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims); extern void oacc_replace_fn_attrib (tree fn, tree dims); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index c9e9dec8559..4c68fca870d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -15164,6 +15164,24 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree fndecl; call_stmt = as_a <gcall *> (stmt); fndecl = gimple_call_fndecl (call_stmt); + if (fndecl + && lookup_attribute ("omp metadirective construct target", + DECL_ATTRIBUTES (fndecl))) + { + bool in_target_ctx = false; + + for (omp_context *up = ctx; up; up = up->outer) + if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET) + { + in_target_ctx = true; + break; + } + if (!ctx || !in_target_ctx) + warning_at (gimple_location (stmt), 0, + "direct calls to an offloadable function containing " + "metadirectives with a %<construct={target}%> " + "selector may produce unexpected results"); + } if (fndecl && fndecl_built_in_p (fndecl, BUILT_IN_NORMAL)) switch (DECL_FUNCTION_CODE (fndecl)) diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 53eeb8b6d96..162cbc58b27 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,8 @@ +2022-01-28 Kwok Cheung Yeung <kcy@codesourcery.com> + + * c-c++-common/gomp/metadirective-4.c (main): Add expected warning. + * gfortran.dg/gomp/metadirective-4.f90 (test): Likewise. + 2022-01-25 Kwok Cheung Yeung <kcy@codesourcery.com> * c-c++-common/gomp/metadirective-7.c: New. diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c index c4b109295db..25efbe046bf 100644 --- a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c @@ -25,7 +25,7 @@ void f(double a[], double x) { /* TODO: This does not execute a version of f with the default clause active as might be expected. */ - f (a, 2.71828); + f (a, 2.71828); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ return 0; } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 index b82c9ea96d9..65eb05cd2fb 100644 --- a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 @@ -13,7 +13,7 @@ program test ! TODO: This does not execute a version of f with the default clause ! active as might be expected. - call f (a, 2.71828) + call f (a, 2.71828) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } contains subroutine f (a, x) integer :: i diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 0d47950d20d..5f4fb9362b5 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,9 @@ +2022-01-28 Kwok Cheung Yeung <kcy@codesourcery.com> + + * testsuite/libgomp.c-c++-common/metadirective-2.c (main): Add + expected warning. + * testsuite/libgomp.fortran/metadirective-2.f90 (test): Likewise. + 2022-01-25 Kwok Cheung Yeung <kcy@codesourcery.com> * Makefile.am (libgomp_la_SOURCES): Add selector.c. diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c index cd5c6c5e21a..55a6098e525 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c @@ -31,7 +31,7 @@ void f(double a[], double x) { /* TODO: This does not execute a version of f with the default clause active as might be expected. */ - f (a, M_E); + f (a, M_E); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ for (i = 0; i < N; i++) if (fabs (a[i] - (M_E * i)) > EPSILON) diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 index 32017a00077..d83474cf2db 100644 --- a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 +++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 @@ -19,7 +19,7 @@ program test ! TODO: This does not execute a version of f with the default clause ! active as might be expected. - call f (a, E_CONST) + call f (a, E_CONST) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } do i = 1, N if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2
reply other threads:[~2022-06-29 14:44 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20220629144409.58049384D15E@sourceware.org \ --to=kcy@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).