From c7c30edd4a6049a0debe0ed17ffdb9c196d4c677 Mon Sep 17 00:00:00 2001 From: Martin Jambor Date: Thu, 8 Feb 2018 14:03:52 +0100 Subject: [PATCH] [hsa] Set program allocation for static local variables 2018-02-08 Martin Jambor * hsa-gen.c (get_symbol_for_decl): Set program allocation for static local variables. libgomp/ * testsuite/libgomp.hsa.c/staticvar.c: New test. From-SVN: r257484 --- gcc/ChangeLog | 5 +++++ gcc/hsa-gen.c | 10 ++++++--- libgomp/ChangeLog | 4 ++++ libgomp/testsuite/libgomp.hsa.c/staticvar.c | 23 +++++++++++++++++++++ 4 files changed, 39 insertions(+), 3 deletions(-) create mode 100644 libgomp/testsuite/libgomp.hsa.c/staticvar.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 4d65b9435a5..770d408fa16 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2018-02-08 Martin Jambor + + * hsa-gen.c (get_symbol_for_decl): Set program allocation for + static local variables. + 2018-02-08 Richard Biener PR tree-optimization/84278 diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index af0b33d658f..55a46b5a16a 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -932,9 +932,13 @@ get_symbol_for_decl (tree decl) else if (lookup_attribute ("hsa_group_segment", DECL_ATTRIBUTES (decl))) segment = BRIG_SEGMENT_GROUP; - else if (TREE_STATIC (decl) - || lookup_attribute ("hsa_global_segment", - DECL_ATTRIBUTES (decl))) + else if (TREE_STATIC (decl)) + { + segment = BRIG_SEGMENT_GLOBAL; + allocation = BRIG_ALLOCATION_PROGRAM; + } + else if (lookup_attribute ("hsa_global_segment", + DECL_ATTRIBUTES (decl))) segment = BRIG_SEGMENT_GLOBAL; else segment = BRIG_SEGMENT_PRIVATE; diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index df8a57d6ab1..d0130490de3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2018-02-08 Martin Jambor + + * testsuite/libgomp.hsa.c/staticvar.c: New test. + 2018-02-07 Rainer Orth * testsuite/libgomp.oacc-c-c++-common/pr84217.c (abort) diff --git a/libgomp/testsuite/libgomp.hsa.c/staticvar.c b/libgomp/testsuite/libgomp.hsa.c/staticvar.c new file mode 100644 index 00000000000..6d20c9aa328 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/staticvar.c @@ -0,0 +1,23 @@ +extern void abort (void); + +#pragma omp declare target +int +foo (void) +{ + static int s; + return ++s; +} +#pragma omp end declare target + +int +main () +{ + int r; + #pragma omp target map(from:r) + { + r = foo (); + } + if (r != 1) + abort (); + return 0; +} -- 2.30.2