diff options
| author | Martin Jambor <mjambor@suse.cz> | 2016-05-16 19:49:44 +0200 |
|---|---|---|
| committer | Martin Jambor <jamborm@gcc.gnu.org> | 2016-05-16 19:49:44 +0200 |
| commit | 51d9ed489006f5693cb3d1bddaba431a98aa86c6 (patch) | |
| tree | 5ed5bcd7d4b7df31c135b193a9f9d59d92b00c1b | |
| parent | 538374e1396abe947817191a4da34ecc0ad3fb19 (diff) | |
| download | gcc-51d9ed489006f5693cb3d1bddaba431a98aa86c6.zip gcc-51d9ed489006f5693cb3d1bddaba431a98aa86c6.tar.gz gcc-51d9ed489006f5693cb3d1bddaba431a98aa86c6.tar.bz2 | |
[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
| -rw-r--r-- | gcc/ChangeLog | 5 | ||||
| -rw-r--r-- | gcc/hsa-gen.c | 19 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 4 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.hsa.c/complex-align-2.c | 27 |
4 files changed, 51 insertions, 4 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 4b3dbb2..5df3ec2 100644 --- 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 --- 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 --- 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 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; +} + |