Commit 2e5efa67 by Cesar Philippidis Committed by Cesar Philippidis

re PR middle-end/84955 (Incorrect OpenACC tile expansion)

PR middle-end/84955

	gcc/
	* lto-streamer-out.c (output_function): Fix CFG loop state before
	streaming out.
	* omp-expand.c (expand_oacc_for): Handle calls to internal
	functions like regular functions.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test.
	* testsuite/libgomp.oacc-fortran/pr84955.f90: New test.

Co-Authored-By: Richard Biener <rguenther@suse.de>

From-SVN: r259346
parent c1566f89
2018-04-12 Cesar Philippidis <cesar@codesourcery.com>
Richard Biener <rguenther@suse.de>
PR middle-end/84955
* lto-streamer-out.c (output_function): Fix CFG loop state before
streaming out.
* omp-expand.c (expand_oacc_for): Handle calls to internal
functions like regular functions.
2018-04-12 Richard Biener <rguenther@suse.de> 2018-04-12 Richard Biener <rguenther@suse.de>
PR lto/85371 PR lto/85371
......
...@@ -2084,6 +2084,9 @@ output_function (struct cgraph_node *node) ...@@ -2084,6 +2084,9 @@ output_function (struct cgraph_node *node)
/* Set current_function_decl and cfun. */ /* Set current_function_decl and cfun. */
push_cfun (fn); push_cfun (fn);
/* Fixup loops if required to match discovery done in the reader. */
loop_optimizer_init (AVOID_CFG_MODIFICATIONS);
/* Make string 0 be a NULL string. */ /* Make string 0 be a NULL string. */
streamer_write_char_stream (ob->string_stream, 0); streamer_write_char_stream (ob->string_stream, 0);
...@@ -2176,12 +2179,13 @@ output_function (struct cgraph_node *node) ...@@ -2176,12 +2179,13 @@ output_function (struct cgraph_node *node)
streamer_write_record_start (ob, LTO_null); streamer_write_record_start (ob, LTO_null);
output_cfg (ob, fn); output_cfg (ob, fn);
pop_cfun ();
} }
else else
streamer_write_uhwi (ob, 0); streamer_write_uhwi (ob, 0);
loop_optimizer_finalize ();
pop_cfun ();
/* Create a section to hold the pickled output of this function. */ /* Create a section to hold the pickled output of this function. */
produce_asm (ob, function); produce_asm (ob, function);
......
...@@ -5439,6 +5439,14 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) ...@@ -5439,6 +5439,14 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd)
split->flags ^= EDGE_FALLTHRU | EDGE_TRUE_VALUE; split->flags ^= EDGE_FALLTHRU | EDGE_TRUE_VALUE;
/* Add a dummy exit for the tiled block when cont_bb is missing. */
if (cont_bb == NULL)
{
edge e = make_edge (body_bb, exit_bb, EDGE_FALSE_VALUE);
e->probability = profile_probability::even ();
split->probability = profile_probability::even ();
}
/* Initialize the user's loop vars. */ /* Initialize the user's loop vars. */
gsi = gsi_start_bb (elem_body_bb); gsi = gsi_start_bb (elem_body_bb);
expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset); expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset);
......
2018-04-12 Cesar Philippidis <cesar@codesourcery.com>
PR middle-end/84955
* testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test.
* testsuite/libgomp.oacc-fortran/pr84955.f90: New test.
2018-04-12 Jakub Jelinek <jakub@redhat.com> 2018-04-12 Jakub Jelinek <jakub@redhat.com>
PR target/85328 PR target/85328
......
/* { dg-do compile } */
int
main ()
{
int i, j;
#pragma acc parallel loop tile(2,3)
for (i = 1; i < 10; i++)
for (j = 1; j < 10; j++)
for (;;)
;
#pragma acc parallel loop
for (i = 1; i < 10; i++)
for (;;)
;
return i + j;
}
! { dg-do compile }
subroutine s
integer :: i, j
!$acc parallel loop tile(2,3)
do i = 1, 10
do j = 1, 10
do
end do
end do
end do
!$acc end parallel loop
!$acc parallel loop
do i = 1, 10
do
end do
end do
!$acc end parallel loop
end subroutine s
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment