[hsa] Increase hsa symbol alignment to natural one

2016-05-16  Martin Jambor  <mjambor@suse.cz>

	* hsa-gen.c (fillup_for_decl): Increase alignment to natural one.
	(get_symbol_for_decl): Sorry if a global symbol in under-aligned.

libgomp/
        * testsuite/libgomp.hsa.c/complex-align-2.c: New test.

From-SVN: r236295

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 4b3dbb2..5df3ec2 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2016-05-16  Martin Jambor  <mjambor@suse.cz>
+
+       * hsa-gen.c (fillup_for_decl): Increase alignment to natural one.
+       (get_symbol_for_decl): Sorry if a global symbol in under-aligned.
+
 2016-05-16  Marek Polacek  <polacek@redhat.com>
 
        * gimple.c (maybe_remove_unused_call_args): Fix typos in the

diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 5baf607..697d599 100644 (file)
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -203,9 +203,13 @@ hsa_symbol::fillup_for_decl (tree decl)
 {
   m_decl = decl;
   m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
-
   if (hsa_seen_error ())
-    m_seen_error = true;
+    {
+      m_seen_error = true;
+      return;
+    }
+
+  m_align = MAX (m_align, hsa_natural_alignment (m_type));
 }
 
 /* Constructor of class representing global HSA function/kernel information and
@@ -929,6 +933,14 @@ get_symbol_for_decl (tree decl)
                                BRIG_LINKAGE_PROGRAM, true,
                                BRIG_ALLOCATION_PROGRAM, align);
          hsa_cfun->m_global_symbols.safe_push (sym);
+         sym->fillup_for_decl (decl);
+         if (sym->m_align > align)
+           {
+             sym->m_seen_error = true;
+             HSA_SORRY_ATV (EXPR_LOCATION (decl),
+                            "HSA specification requires that %E is at least "
+                            "naturally aligned", decl);
+           }
        }
       else
        {
@@ -944,12 +956,11 @@ get_symbol_for_decl (tree decl)
          sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_PRIVATE,
                                BRIG_LINKAGE_FUNCTION);
          sym->m_align = align;
+         sym->fillup_for_decl (decl);
          hsa_cfun->m_private_variables.safe_push (sym);
        }
 
-      sym->fillup_for_decl (decl);
       sym->m_name = hsa_get_declaration_name (decl);
-
       *slot = sym;
       return sym;
     }

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 9de04f5..f509114 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2016-05-16  Martin Jambor  <mjambor@suse.cz>
+
+        * testsuite/libgomp.hsa.c/complex-align-2.c: New test.
+
 2016-05-02  Nathan Sidwell  <nathan@codesourcery.com>
 
        * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust

diff --git a/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c
new file mode 100644 (file)
index 0000000..b2d7acf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c
@@ -0,0 +1,27 @@
+#pragma omp declare target
+    _Complex int *g;
+#pragma omp end declare target
+
+
+
+_Complex float f(void);
+
+int
+main ()
+{
+  _Complex int y;
+#pragma omp target map(from:y)
+  {
+    _Complex int x;
+    g = &x;
+    __imag__ x = 1;
+    __real__ x = 2;
+    y = x;
+  }
+
+  if ((__imag__ y != 1)
+      || (__real__ y != 2))
+    __builtin_abort ();
+  return 0;
+}
+