[PATCH] openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly

[Kwok Cheung Yeung <kcy@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 3035 bytes --]

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
> 

[-- Attachment #2: 0001-openmp-Add-warning-when-functions-containing-metadir.patch --]
[-- Type: text/plain, Size: 8613 bytes --]

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(-)

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