diff options
| author | Nathan Sidwell <nathan@codesourcery.com> | 2016-08-03 17:26:51 +0000 |
|---|---|---|
| committer | Nathan Sidwell <nathan@gcc.gnu.org> | 2016-08-03 17:26:51 +0000 |
| commit | c1311c86c0c1057b43146c8817dfe453c7a88fd4 (patch) | |
| tree | 29f0919ef5bd3f0b99f39f7f69700e3f71b2f081 | |
| parent | 0bd9bdb4dbc21dbdef0955db2645df9e97f83c82 (diff) | |
| download | gcc-c1311c86c0c1057b43146c8817dfe453c7a88fd4.zip gcc-c1311c86c0c1057b43146c8817dfe453c7a88fd4.tar.gz gcc-c1311c86c0c1057b43146c8817dfe453c7a88fd4.tar.bz2 | |
nvptx.c (nvptx_declare_function_name): Round frame size to DImode boundary.
gcc/
* config/nvptx/nvptx.c (nvptx_declare_function_name): Round frame
size to DImode boundary.
(nvptx_propagate): Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
From-SVN: r239086
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 12 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 4 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c | 28 |
4 files changed, 46 insertions, 4 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 3da9eb7..2223290 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2016-08-03 Nathan Sidwell <nathan@codesourcery.com> + + * config/nvptx/nvptx.c (nvptx_declare_function_name): Round frame + size to DImode boundary. + (nvptx_propagate): Likewise. + 2016-08-03 Alan Modra <amodra@gmail.com> * config/rs6000/rs6000.h (SLOW_UNALIGNED_ACCESS): Make scalar diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 6c78699..4d87ead 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -1037,11 +1037,14 @@ nvptx_declare_function_name (FILE *file, const char *name, const_tree decl) init_frame (file, STACK_POINTER_REGNUM, UNITS_PER_WORD, crtl->outgoing_args_size); - /* Declare a local variable for the frame. */ + /* Declare a local variable for the frame. Force its size to be + DImode-compatible. */ HOST_WIDE_INT sz = get_frame_size (); if (sz || cfun->machine->has_chain) init_frame (file, FRAME_POINTER_REGNUM, - crtl->stack_alignment_needed / BITS_PER_UNIT, sz); + crtl->stack_alignment_needed / BITS_PER_UNIT, + (sz + GET_MODE_SIZE (DImode) - 1) + & ~(HOST_WIDE_INT)(GET_MODE_SIZE (DImode) - 1)); /* Declare the pseudos we have as ptx registers. */ int maxregs = max_reg_num (); @@ -3266,8 +3269,9 @@ nvptx_propagate (basic_block block, rtx_insn *insn, propagate_mask rw, rtx pred = NULL_RTX; rtx_code_label *label = NULL; - gcc_assert (!(fs & (GET_MODE_SIZE (DImode) - 1))); - fs /= GET_MODE_SIZE (DImode); + /* The frame size might not be DImode compatible, but the frame + array's declaration will be. So it's ok to round up here. */ + fs = (fs + GET_MODE_SIZE (DImode) - 1) / GET_MODE_SIZE (DImode); /* Detect single iteration loop. */ if (fs == 1) fs = 0; diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ebb4bc9..850188f 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2016-08-03 Nathan Sidwell <nathan@codesourcery.com> + + * testsuite/libgomp.oacc-c-c++-common/crash-1.c: New. + 2016-07-15 Cesar Philippidis <cesar@codesourcery.com> * testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c new file mode 100644 index 0000000..a75a817 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c @@ -0,0 +1,28 @@ +/* { dg-do compile } */ +/* { dg-options "-O0" } */ + +/* ICEd in nvptx backend due to unexpected frame size. */ +#pragma acc routine worker +void +worker_matmul (int *c, int i) +{ + int j; + +#pragma acc loop + for (j = 0; j < 4; j++) + c[j] = j; +} + + +int +main () +{ + int c[4]; + +#pragma acc parallel + { + worker_matmul (c, 0); + } + + return 0; +} |