From 51d9ed489006f5693cb3d1bddaba431a98aa86c6 Mon Sep 17 00:00:00 2001 From: Martin Jambor Date: Mon, 16 May 2016 19:49:44 +0200 Subject: [PATCH] [hsa] Increase hsa symbol alignment to natural one 2016-05-16 Martin Jambor * 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 --- gcc/ChangeLog | 5 ++++ gcc/hsa-gen.c | 19 ++++++++++--- libgomp/ChangeLog | 4 +++ .../testsuite/libgomp.hsa.c/complex-align-2.c | 27 +++++++++++++++++++ 4 files changed, 51 insertions(+), 4 deletions(-) create mode 100644 libgomp/testsuite/libgomp.hsa.c/complex-align-2.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 4b3dbb23f51..5df3ec2da83 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2016-05-16 Martin Jambor + + * 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 * gimple.c (maybe_remove_unused_call_args): Fix typos in the diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index 5baf6073e3b..697d5997519 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 9de04f57d73..f509114a20c 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2016-05-16 Martin Jambor + + * testsuite/libgomp.hsa.c/complex-align-2.c: New test. + 2016-05-02 Nathan Sidwell * 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 00000000000..b2d7acff443 --- /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.30.2