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 <k...@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 <k...@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(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 78bae567ae4..c8a01a4ca52 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -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);
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index 4edeb58cc73..842e9fd868f 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -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. */
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index fdde4a3dfb0..ccdea015e15 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 59804300c28..07613362ef0 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -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))
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/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
--
2.25.1
