Hi,
in the last round fo alignment fixes, we have forgot to make sure that
all symbols are at least naturally aligned, which is a hard HSAIL
requirement. This caused problems when emitting a symbol for a
private complex number, as the natural alignment as defined by HSAIL
is twice the one of the component, which was selected by gcc.
The following patch addresses this in two ways. First, it simply
increases the alignment of symbols that are only accessible from
within HSAIL. If however a symbol that is shared in between host and
an HSA accelerator is under-aligned (in my experience it only happens
if the user uses the aligned attribute), we have no option but to
abort HSAIL generation because even if we did generate it, it would
not finalize.
Bootstrapped and tested on x86_64-linux with hsa enabled. I will
commit it to trunk and the gcc-6 branch shortly.
Martin
2016-05-16 Martin Jambor <mjam...@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.
---
gcc/hsa-gen.c | 19 ++++++++++++----
libgomp/testsuite/libgomp.hsa.c/complex-align-2.c | 27 +++++++++++++++++++++++
2 files changed, 42 insertions(+), 4 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.hsa.c/complex-align-2.c
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 5baf607..697d599 100644
--- 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/testsuite/libgomp.hsa.c/complex-align-2.c
b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c
new file mode 100644
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;
+}
+
--
2.8.2
