public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] openmp: Add warning when functions containing metadirectives with 'construct={target}' called direct
@ 2022-06-29 14:44 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:44 UTC (permalink / raw)
To: gcc-cvs
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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:44 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:44 [gcc/devel/omp/gcc-12] openmp: Add warning when functions containing metadirectives with 'construct={target}' called direct Kwok Yeung
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).