@@ -14775,6 +14775,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);
@@ -2922,6 +2922,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. */
@@ -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);
@@ -14300,6 +14300,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))
@@ -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;
}
@@ -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
@@ -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)
@@ -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
Hello Regarding this issue which we discussed previously - I have created a patch that adds a warning when this situation is detected. When a metadirective in a explicitly marked target function is gimplified, it is checked to see if it contains a 'construct={target}' selector - if it does, then the containing function is marked with 'omp metadirective construct target'. In the omp-low pass, when function calls are processed, the target function is checked to see if it contains the marker. If it does and the call is not made in a target context, a warning is emitted. This will obviously not catch every possible occurence (e.g. if the function containing the metadirective is called from another target function which is then called locally, or if the call is made via a function pointer), but it might still be useful? Okay for mainline (once the metadirective patches are done)? Thanks Kwok On 26/07/2021 10:23 pm, Jakub Jelinek wrote: > On Mon, Jul 26, 2021 at 10:19:35PM +0100, Kwok Cheung Yeung wrote: >>> Yes, that is a target variant, but I'm pretty sure we've decided that >>> the target construct added for declare target is actually not a dynamic >>> property. So basically mostly return to the 5.0 wording with clarifications >>> for Fortran. See >>> https://github.com/OpenMP/spec/issues/2612#issuecomment-849742988 >>> for details. >>> Making the target in construct dynamic would pretty much force all the >>> scoring to be dynamic as well. >> >> In that comment, Deepak says: >> >> So, we decided to keep the target trait static, requiring that the declare >> target directive must be explicit and that the function version must be >> different from the version of the function that may be called outside of a >> target region (with the additional clarification that whether it differs or >> not will be implementation defined). >> >> "the function version must be different from the version of the function >> that may be called outside of a target region": This is what we do not have >> in GCC at the moment - the function versions called within and outside >> target regions are the same on the host. >> >> "whether it differs or not will be implementation defined": So whether a >> function with 'declare target' and a metadirective involving a 'target' >> construct behaves the same or not when called from both inside and outside >> of a target region is implementation defined? >> >> I will leave the treatment of target constructs in the selector as it is >> then, with both calls going to the same function with the metadirective >> resolving to the 'target' variant. I will try to address your other concerns >> later. > > I think you're right, it should differ in the host vs. target version iff > it is in explicit declare target block, my memory is weak, but let's implement > the 5.0 wording for now (and ignore the 5.1 wording later on) and only when > we'll be doing 5.2 change this (and change for both metadirective and > declare variant at that point). > Ok? > > Jakub > From 741b037a8cd6b85d43a6273ab305ce07705dfa23 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung <kcy@codesourcery.com> Date: Fri, 28 Jan 2022 13:56:33 +0000 Subject: [PATCH] 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. --- gcc/gimplify.cc | 21 +++++++++++++++++++ gcc/omp-general.cc | 21 +++++++++++++++++++ gcc/omp-general.h | 1 + gcc/omp-low.cc | 18 ++++++++++++++++ .../c-c++-common/gomp/metadirective-4.c | 2 +- .../gfortran.dg/gomp/metadirective-4.f90 | 2 +- .../libgomp.c-c++-common/metadirective-2.c | 2 +- .../libgomp.fortran/metadirective-2.f90 | 2 +- 8 files changed, 65 insertions(+), 4 deletions(-)