openmp: Add warning when functions containing metadirectives with 'construct={target...

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 --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index eb3a1130f727c922bcfde83ba9bf621b9a4fe4bf..abcdb0203576948a81c5e904cc57fa54acc1dbc0 100644 (file)
--- 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 fca7f4d6bf225647d04980dbfbfbc62eeafbf78b..105bd81b5b8b3d558e9d2c71e515de77dd186246 100644 (file)
--- 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 48953e11997fdd4c7fb1a54675cef61c186dc333..05b960ea983019c0df40308c02b4ac6275f0449c 100644 (file)
--- 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 65c1a60ad41927c375dd2e0d220a5f28af8ef560..6074b0214f412a6a43a6b1b6586a51b577114326 100644 (file)
--- 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 c9e9dec85595b130f88776c9203595f3f65129d9..4c68fca870dade430c218c5eaed094fe85cb98c7 100644 (file)
--- 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 53eeb8b6d9655a56a0a65baff6307e97fed0c3a7..162cbc58b2778dad6301d7ad1d8a693f0487dca9 100644 (file)
--- 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 c4b109295db25e57b82de977a4bbe707e78477dc..25efbe046bf23367edba6cac17b19f0a8fdc678d 100644 (file)
--- 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 b82c9ea96d96589a2dad7b90d0bb34fced7b13b7..65eb05cd2fbaf6d623d3a8fc5d619475426b07fc 100644 (file)
--- 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 0d47950d20d24bad14f426ea2d662f2291e44bad..5f4fb9362b596a62ed635abd3188f5756026206c 100644 (file)
--- 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 cd5c6c5e21a16522ac819cbd98ec406155a8aae4..55a6098e525a0ca4d8bd282d75712ea964da963f 100644 (file)
--- 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 32017a000779fafc2009b0bdcf11bf126d3daf4f..d83474cf2db991d25f1c70849dbbac3fe84205fd 100644 (file)
--- 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