This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[hsa] Set program allocation for static local variables
- From: Martin Jambor <mjambor at suse dot cz>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc:
- Date: Thu, 08 Feb 2018 13:57:22 +0100
- Subject: [hsa] Set program allocation for static local variables
- Authentication-results: sourceware.org; auth=none
Hi,
it has been brought to my attention that libgomp.c/target-28.c
testcase fails to finalize because the static variable s has illegal
hsa allocation. Fixed by the patch below, which I am about to commit
to trunk (and will commit to the gcc-7-branch after testing there).
Thanks,
Martin
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.
Added testcase
---
gcc/hsa-gen.c | 10 +++++++---
libgomp/testsuite/libgomp.hsa.c/staticvar.c | 23 +++++++++++++++++++++++
2 files changed, 30 insertions(+), 3 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.hsa.c/staticvar.c
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/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.15.1