aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMartin Jambor <mjambor@suse.cz>2018-02-08 14:03:52 +0100
committerMartin Jambor <jamborm@gcc.gnu.org>2018-02-08 14:03:52 +0100
commitc7c30edd4a6049a0debe0ed17ffdb9c196d4c677 (patch)
tree1ddc2a0a9a52e7b9fbe52d94618bbcaeb647280b
parent414fef4e668856edb8ca885525679c5d5e691fd1 (diff)
downloadgcc-c7c30edd4a6049a0debe0ed17ffdb9c196d4c677.zip
gcc-c7c30edd4a6049a0debe0ed17ffdb9c196d4c677.tar.gz
gcc-c7c30edd4a6049a0debe0ed17ffdb9c196d4c677.tar.bz2
[hsa] Set program allocation for static local variables
2018-02-08 Martin Jambor <mjambor@suse.cz> * 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
-rw-r--r--gcc/ChangeLog5
-rw-r--r--gcc/hsa-gen.c10
-rw-r--r--libgomp/ChangeLog4
-rw-r--r--libgomp/testsuite/libgomp.hsa.c/staticvar.c23
4 files changed, 39 insertions, 3 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 4d65b94..770d408 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2018-02-08 Martin Jambor <mjambor@suse.cz>
+
+ * hsa-gen.c (get_symbol_for_decl): Set program allocation for
+ static local variables.
+
2018-02-08 Richard Biener <rguenther@suse.de>
PR tree-optimization/84278
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index af0b33d..55a46b5 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 df8a57d..d013049 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2018-02-08 Martin Jambor <mjambor@suse.cz>
+
+ * testsuite/libgomp.hsa.c/staticvar.c: New test.
+
2018-02-07 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE>
* 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 0000000..6d20c9a
--- /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;
+}