LIV2/bebbo-gcc
1
0
Fork 0
mirror of https://github.com/LIV2/bebbo-gcc.git synced 2025-12-06 00:23:32 +00:00
Code Issues Packages Projects Releases Wiki Activity

[PTX] fix worker propagation ICE

Browse Source 
Backport trunk r239086:

	gcc/
	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.

	libgomp/
	2016-08-03  Nathan Sidwell  <nathan@codesourcery.com>

	* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.

Backport trunk r239125 'Make libgomp.oacc-c-c++-common/crash-1.c a "link" test,
and don't hardcode -O0':

	libgomp/
	2016-08-04  Thomas Schwinge  <thomas@codesourcery.com>

	* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
	test, and don't hardcode -O0.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-6-branch@244261 138bc75d-0d04-0410-961f-82ee72b054a4
This commit is contained in:
tschwinge 2017-01-10 12:07:13 +00:00
parent 8682ea4d79
commit 71f5a3c500
4 changed files with 57 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
gcc/ChangeLog
View File

@ -1,3 +1,12 @@
2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
Backport trunk r239086:
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.
2017-01-10 Chung-Ju Wu <jasonwucj@gmail.com>
Backport from mainline

12
gcc/config/nvptx/nvptx.c
View File

@ -989,11 +989,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 ();
@ -3212,8 +3215,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;

13
libgomp/ChangeLog
View File

@ -1,3 +1,16 @@
2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
Backport trunk r239125:
2016-08-04 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
test, and don't hardcode -O0.
Backport trunk r239086:
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
2016-12-21 Release Manager
* GCC 6.3.0 released.

27
libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c Normal file
View File

@ -0,0 +1,27 @@
/* { dg-do link } */
/* For -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;
}