mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-15 09:00:29 +08:00
Code Issues Packages Projects Releases Wiki Activity

OpenMP: Fix 'omp declare target' handling for vars [PR99509]

Browse Source 
For variables with 'declare target' attribute,
varpool_node::get_create marks variables as offload; however,
if the node already exists, it is not updated. C/C++ may tag
decl with 'declare target implicit', which may only be after
varpool creation turned into 'declare target' or 'declare target link';
in this case, the tagging has to happen in the FE.

gcc/c/ChangeLog:

	PR c++/99509
	* c-decl.c (finish_decl): For 'omp declare target implicit' vars,
	ensure that the varpool node is marked as offloadable.

gcc/cp/ChangeLog:

	PR c++/99509
	* decl.c (cp_finish_decl): For 'omp declare target implicit' vars,
	ensure that the varpool node is marked as offloadable.

libgomp/ChangeLog:

	PR c++/99509
	* testsuite/libgomp.c-c++-common/declare_target-1.c: New test.
This commit is contained in:
Tobias Burnus 2021-03-15 10:12:58 +01:00
parent b516a15371
commit f20fe2cb21
3 changed files with 59 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
gcc/c/c-decl.c
View File

@ -58,6 +58,9 @@ along with GCC; see the file COPYING3. If not see
#include "c-family/name-hint.h"
#include "c-family/known-headers.h"
#include "c-family/c-spellcheck.h"
#include "context.h" /* For 'g'. */
#include "omp-general.h"
#include "omp-offload.h" /* For offload_vars. */
#include "tree-pretty-print.h"
@ -5658,9 +5661,22 @@ finish_decl (tree decl, location_t init_loc, tree init,
DECL_ATTRIBUTES (decl))
&& !lookup_attribute ("omp declare target link",
DECL_ATTRIBUTES (decl)))
DECL_ATTRIBUTES (decl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (decl));
{
DECL_ATTRIBUTES (decl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (decl));
symtab_node *node = symtab_node::get (decl);
if (node != NULL)
{
node->offloadable = 1;
if (ENABLE_OFFLOADING)
{
g->have_offload = true;
if (is_a <varpool_node *> (node))
vec_safe_push (offload_vars, decl);
}
}
}
}
/* This is the last point we can lower alignment so give the target the

21
gcc/cp/decl.c
View File

@ -53,7 +53,9 @@ along with GCC; see the file COPYING3. If not see
#include "asan.h"
#include "gcc-rich-location.h"
#include "langhooks.h"
#include "context.h" /* For 'g'. */
#include "omp-general.h"
#include "omp-offload.h" /* For offload_vars. */
/* Possible cases of bad specifiers type used by bad_specifiers. */
enum bad_spec_place {
@ -8176,9 +8178,22 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p,
DECL_ATTRIBUTES (decl))
&& !lookup_attribute ("omp declare target link",
DECL_ATTRIBUTES (decl)))
DECL_ATTRIBUTES (decl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (decl));
{
DECL_ATTRIBUTES (decl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (decl));
symtab_node *node = symtab_node::get (decl);
if (node != NULL)
{
node->offloadable = 1;
if (ENABLE_OFFLOADING)
{
g->have_offload = true;
if (is_a <varpool_node *> (node))
vec_safe_push (offload_vars, decl);
}
}
}
}
/* This is the last point we can lower alignment so give the target the

22
libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c Normal file
View File

@ -0,0 +1,22 @@
/* PR c++/99509 */
#pragma omp declare target
int data[] = {5};
#pragma omp end declare target
static inline int
foo (int idx)
{
return data[idx];
}
int
main ()
{
int i = -1;
#pragma omp target map(from:i)
i = foo(0);
if (i != 5)
__builtin_abort ();
return 0;
}