From: Martin Jambor <mjambor@suse.cz>
Date: Thu, 8 Feb 2018 13:03:52 +0000 (+0100)
Subject: [hsa] Set program allocation for static local variables
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=c7c30edd4a6049a0debe0ed17ffdb9c196d4c677;p=gcc.git

[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
---

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  <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 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  <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 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;
+}