OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

openmp: Also implicitly mark as declare target to functions mentioned in target regions 

OpenMP 5.0 also specifies that functions referenced from target regions
(except for target regions with device(ancestor:)) are also implicitly declare target to.

This patch implements that.

2020-05-14  Jakub Jelinek  <jakub@redhat.com>

	* function.h (struct function): Add has_omp_target bit.
	* omp-offload.c (omp_discover_declare_target_fn_r): New function,
	old renamed to ...
	(omp_discover_declare_target_tgt_fn_r): ... this.
	(omp_discover_declare_target_var_r): Call
	omp_discover_declare_target_tgt_fn_r instead of
	omp_discover_declare_target_fn_r.
	(omp_discover_implicit_declare_target): Also queue functions with
	has_omp_target bit set, for those walk with
	omp_discover_declare_target_fn_r, for declare target to functions
	walk with omp_discover_declare_target_tgt_fn_r.
gcc/c/
	* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
gcc/cp/
	* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
gcc/fortran/
	* trans-openmp.c: Include function.h.
	(gfc_trans_omp_target): Set cfun->has_omp_target.
libgomp/
	* testsuite/libgomp.c-c++-common/target-40.c: New test.
This commit is contained in:
Jakub Jelinek 2020-05-14 09:48:32 +02:00
parent 42ef8a5e66
commit 49ddde69fc
11 changed files with 129 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
gcc/ChangeLog
View File

@ -1,3 +1,17 @@
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* function.h (struct function): Add has_omp_target bit.
* omp-offload.c (omp_discover_declare_target_fn_r): New function,
old renamed to ...
(omp_discover_declare_target_tgt_fn_r): ... this.
(omp_discover_declare_target_var_r): Call
omp_discover_declare_target_tgt_fn_r instead of
omp_discover_declare_target_fn_r.
(omp_discover_implicit_declare_target): Also queue functions with
has_omp_target bit set, for those walk with
omp_discover_declare_target_fn_r, for declare target to functions
walk with omp_discover_declare_target_tgt_fn_r.
2020-05-14 Uroš Bizjak <ubizjak@gmail.com>
PR target/95046

4
gcc/c/ChangeLog
View File

@ -1,3 +1,7 @@
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
2020-05-07 Richard Biener <rguenther@suse.de>
PR middle-end/94703

1
gcc/c/c-parser.c
View File

@ -19866,6 +19866,7 @@ check_clauses:
}
pc = &OMP_CLAUSE_CHAIN (*pc);
}
cfun->has_omp_target = true;
return true;
}

4
gcc/cp/ChangeLog
View File

@ -1,3 +1,7 @@
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
2020-05-13 Patrick Palka <ppalka@redhat.com>
PR c++/79706

4
gcc/cp/cp-gimplify.c
View File

@ -1558,6 +1558,10 @@ cp_genericize_r (tree *stmt_p, int *walk_subtrees, void *data)
}
break;
case OMP_TARGET:
cfun->has_omp_target = true;
break;
case TRY_BLOCK:
{
*walk_subtrees = 0;

5
gcc/fortran/ChangeLog
View File

@ -1,3 +1,8 @@
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* trans-openmp.c: Include function.h.
(gfc_trans_omp_target): Set cfun->has_omp_target.
2020-05-13 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/93497

2
gcc/fortran/trans-openmp.c
View File

@ -44,6 +44,7 @@ along with GCC; see the file COPYING3. If not see
#undef GCC_DIAG_STYLE
#define GCC_DIAG_STYLE __gcc_gfc__
#include "attribs.h"
#include "function.h"
int ompws_flags;
@ -5392,6 +5393,7 @@ gfc_trans_omp_target (gfc_code *code)
omp_clauses);
if (code->op != EXEC_OMP_TARGET)
OMP_TARGET_COMBINED (stmt) = 1;
cfun->has_omp_target = true;
}
gfc_add_expr_to_block (&block, stmt);
return gfc_finish_block (&block);

3
gcc/function.h
View File

@ -421,6 +421,9 @@ struct GTY(()) function {
/* Set if this is a coroutine-related function. */
unsigned int coroutine_component : 1;
/* Set if there are any OMP_TARGET regions in the function. */
unsigned int has_omp_target : 1;
};
/* Add the decl D to the local_decls list of FUN. */

46
gcc/omp-offload.c
View File

@ -190,7 +190,7 @@ omp_declare_target_var_p (tree decl)
declare target to. */
static tree
omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
{
if (TREE_CODE (*tp) == FUNCTION_DECL
&& !omp_declare_target_fn_p (*tp)
@ -219,6 +219,24 @@ omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
return NULL_TREE;
}
/* Similarly, but ignore references outside of OMP_TARGET regions. */
static tree
omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
{
if (TREE_CODE (*tp) == OMP_TARGET)
{
/* And not OMP_DEVICE_ANCESTOR. */
walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
omp_discover_declare_target_tgt_fn_r,
data);
*walk_subtrees = 0;
}
else if (TYPE_P (*tp))
*walk_subtrees = 0;
return NULL_TREE;
}
/* Helper function for omp_discover_implicit_declare_target, called through
walk_tree. Mark referenced FUNCTION_DECLs implicitly as
declare target to. */
@ -227,7 +245,7 @@ static tree
omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
{
if (TREE_CODE (*tp) == FUNCTION_DECL)
return omp_discover_declare_target_fn_r (tp, walk_subtrees, data);
return omp_discover_declare_target_tgt_fn_r (tp, walk_subtrees, data);
else if (VAR_P (*tp)
&& is_global_var (*tp)
&& !omp_declare_target_var_p (*tp))
@ -271,22 +289,32 @@ omp_discover_implicit_declare_target (void)
auto_vec<tree> worklist;
FOR_EACH_DEFINED_FUNCTION (node)
if (omp_declare_target_fn_p (node->decl) && DECL_SAVED_TREE (node->decl))
worklist.safe_push (node->decl);
if (DECL_SAVED_TREE (node->decl))
{
if (omp_declare_target_fn_p (node->decl))
worklist.safe_push (node->decl);
else if (DECL_STRUCT_FUNCTION (node->decl)
&& DECL_STRUCT_FUNCTION (node->decl)->has_omp_target)
worklist.safe_push (node->decl);
}
FOR_EACH_STATIC_INITIALIZER (vnode)
if (omp_declare_target_var_p (vnode->decl))
worklist.safe_push (vnode->decl);
while (!worklist.is_empty ())
{
tree decl = worklist.pop ();
if (TREE_CODE (decl) == FUNCTION_DECL)
walk_tree_without_duplicates (&DECL_SAVED_TREE (decl),
omp_discover_declare_target_fn_r,
&worklist);
else
if (VAR_P (decl))
walk_tree_without_duplicates (&DECL_INITIAL (decl),
omp_discover_declare_target_var_r,
&worklist);
else if (omp_declare_target_fn_p (decl))
walk_tree_without_duplicates (&DECL_SAVED_TREE (decl),
omp_discover_declare_target_tgt_fn_r,
&worklist);
else
walk_tree_without_duplicates (&DECL_SAVED_TREE (decl),
omp_discover_declare_target_fn_r,
&worklist);
}
}

4
libgomp/ChangeLog
View File

@ -1,3 +1,7 @@
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c-c++-common/target-40.c: New test.
2020-05-13 Tobias Burnus <tobias@codesourcery.com>
PR fortran/94690

51
libgomp/testsuite/libgomp.c-c++-common/target-40.c Normal file
View File

@ -0,0 +1,51 @@
/* { dg-do run } */
/* { dg-options "-O0" } */
extern
#ifdef __cplusplus
"C"
#endif
void abort (void);
volatile int v;
#pragma omp declare target to (v)
typedef void (*fnp1) (void);
typedef fnp1 (*fnp2) (void);
void f1 (void) { v++; }
void f2 (void) { v += 4; }
void f3 (void) { v += 16; f1 (); }
fnp1 f4 (void) { v += 64; return f2; }
int a = 1;
int *b = &a;
int **c = &b;
fnp2 f5 (void) { f3 (); return f4; }
#pragma omp declare target to (c)
int
main ()
{
int err = 0;
#pragma omp target map(from:err)
{
volatile int xa;
int *volatile xb;
int **volatile xc;
fnp2 xd;
fnp1 xe;
err = 0;
xa = a;
err |= xa != 1;
xb = b;
err |= xb != &a;
xc = c;
err |= xc != &b;
xd = f5 ();
err |= v != 17;
xe = xd ();
err |= v != 81;
xe ();
err |= v != 85;
}
if (err)
abort ();
return 0;
}