[gomp4, committed] Handle nested loops in kernels regions

[gomp4, committed] Handle nested loops in kernels regions

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1578 bytes --]

Hi,

I.

This patch allows parallelization of an outer loop in an openacc kernels 
region.

The testcase is based on autopar/outer-1.c.


II.

We rely on pass_lim to move the *.omp_data_i loads out of the loop nest. 
For the test-case, pass_lim was managing to move the load out of the 
inner loop, but not the outer loop, because the load was classified as 
'MOVE_PRESERVE_EXECUTION'. By marking the *.omp_data_i load 
non-trapping, it's now classified as 'MOVE_POSSIBLE', and moved out of 
the loop nest.


III.

The 'loops_state_set (LOOPS_NEED_FIXUP)' is a somewhat blunt and 
temporary fix for the oacc kernels variant of PR66846 - parloops does 
not always mark loops for fixup if needed.

The original PR needs an added verify_loop_structure to trigger the 
problem. Normally the problem is hidden by the fact that the first pass 
that runs on the new function is pass_fixup_cfg, which happens to fixup 
the loops (The loops are fixed up because TODO_cleanup_cfg is set during 
pass_fixup_cfg, because the function contains an ECF_CONST function: 
__builtin_omp_get_num_threads).

For the oacc kernels variant, the problem triggers without adding 
verify_loop_structure. During pass_ipa_inline, we call 
loop_optimizer_init, which (given that LOOPS_NEED_FIXUP is not set) 
verifies the loop structure, which fails. Pass_fixup_cfg is not run 
inbetween the discovery of the new function and pass_ipa_inline.


IV.

I've committed this patch to gomp-4_0-branch.

Bootstrapped and reg-tested on x86_64. Build and reg-tested on setup 
with nvidia accelerator.

Thanks,
- Tom

[-- Attachment #2: 0001-Handle-nested-loops-in-kernels-regions.patch --]
[-- Type: text/x-patch, Size: 4374 bytes --]

Handle nested loops in kernels regions

2015-07-12  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (build_receiver_ref): Mark *.omp_data_i as non-trapping.
	* tree-parloops.c (gen_parallel_loop): Add LOOPS_NEED_FIXUP to loop
	state.
	(parallelize_loops): Allow nested loops.

	* testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c: New test.

	* c-c++-common/goacc/kernels-loop-nest.c: New test.
---
 gcc/omp-low.c                                      |  1 +
 .../c-c++-common/goacc/kernels-loop-nest.c         | 42 ++++++++++++++++++++++
 gcc/tree-parloops.c                                |  5 +--
 .../libgomp.oacc-c-c++-common/kernels-loop-nest.c  | 26 ++++++++++++++
 4 files changed, 70 insertions(+), 4 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 11ac909..a938ce0 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1147,6 +1147,7 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
     field = x;
 
   x = build_simple_mem_ref (ctx->receiver_decl);
+  TREE_THIS_NOTRAP (x) = 1;
   x = omp_build_component_ref (x, field);
   if (by_ref)
     x = build_simple_mem_ref (x);
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
new file mode 100644
index 0000000..3e06c9f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
@@ -0,0 +1,42 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* Based on autopar/outer-1.c.  */
+
+#include <stdlib.h>
+
+#define N 1000
+
+int
+main (void)
+{
+  int x[N][N];
+
+#pragma acc kernels copyout (x)
+  {
+    for (int ii = 0; ii < N; ii++)
+      for (int jj = 0; jj < N; jj++)
+	x[ii][jj] = ii + jj + 3;
+  }
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      if (x[i][j] != i + j + 3)
+	abort ();
+
+  return 0;
+}
+
+/* Check that only one loop is analyzed, and that it can be parallelized.  */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* Check that the loop has been split off into a function.  */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
+/* { dg-final { cleanup-tree-dump "optimized" } } */
diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c
index 04708c0..492ffcb 100644
--- a/gcc/tree-parloops.c
+++ b/gcc/tree-parloops.c
@@ -2442,6 +2442,7 @@ gen_parallel_loop (struct loop *loop,
   /* Cancel the loop (it is simpler to do it here rather than to teach the
      expander to do it).  */
   cancel_loop_tree (loop);
+  loops_state_set (LOOPS_NEED_FIXUP);
 
   /* Free loop bound estimations that could contain references to
      removed statements.  */
@@ -2761,10 +2762,6 @@ parallelize_loops (bool oacc_kernels_p)
 	  if (!loop->in_oacc_kernels_region)
 	    continue;
 
-	  /* TODO: Allow nested loops.  */
-	  if (loop->inner)
-	    continue;
-
 	  if (dump_file && (dump_flags & TDF_DETAILS))
 	    fprintf (dump_file,
 		     "Trying loop %d with header bb %d in oacc kernels region\n",
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c
new file mode 100644
index 0000000..21d2599
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 1000
+
+int
+main (void)
+{
+  int x[N][N];
+
+#pragma acc kernels copyout (x)
+  {
+    for (int ii = 0; ii < N; ii++)
+      for (int jj = 0; jj < N; jj++)
+	x[ii][jj] = ii + jj + 3;
+  }
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      if (x[i][j] != i + j + 3)
+	abort ();
+
+  return 0;
+}
-- 
1.9.1

Re: [gomp4, committed] Handle nested loops in kernels regions

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 3541 bytes --]

Hi!

On Sun, 12 Jul 2015 14:46:02 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> This patch allows parallelization of an outer loop in an openacc kernels 
> region.

\o/


> We rely on pass_lim to move the *.omp_data_i loads out of the loop nest. 
> For the test-case, pass_lim was managing to move the load out of the 
> inner loop, but not the outer loop, because the load was classified as 
> 'MOVE_PRESERVE_EXECUTION'. By marking the *.omp_data_i load 
> non-trapping, it's now classified as 'MOVE_POSSIBLE', and moved out of 
> the loop nest.

Should this go into trunk already?  (Jakub?)  Do we need to audit the
code for constructs that need similar treatment?

	* omp-low.c (build_receiver_ref): Mark *.omp_data_i as non-trapping.

--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1147,6 +1147,7 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
     field = x;
 
   x = build_simple_mem_ref (ctx->receiver_decl);
+  TREE_THIS_NOTRAP (x) = 1;
   x = omp_build_component_ref (x, field);
   if (by_ref)
     x = build_simple_mem_ref (x);


> I've committed this patch to gomp-4_0-branch.

Are you planning to also add a Fortran test case?


> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
> @@ -0,0 +1,42 @@
> +/* { dg-additional-options "-O2" } */
> +/* { dg-additional-options "-ftree-parallelize-loops=32" } */
> +/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
> +/* { dg-additional-options "-fdump-tree-optimized" } */
> +[...]
> +/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
> +/* { dg-final { cleanup-tree-dump "optimized" } } */

Committed in r225728:

commit c2da6458c51cc71dccec2e49481560b91d57aa1c
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Mon Jul 13 08:14:05 2015 +0000

    cleanup-tree-dump is no more
    
        ERROR: (DejaGnu) proc "cleanup-tree-dump parloops_oacc_kernels" does not exist.
    
    	gcc/testsuite/
    	* c-c++-common/goacc/kernels-loop-nest.c: Remove cleanup-tree-dump
    	directives.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@225728 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog.gomp                         |    5 +++++
 gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c |    3 ---
 2 files changed, 5 insertions(+), 3 deletions(-)

diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 636a087..4694a60 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-07-13  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/kernels-loop-nest.c: Remove cleanup-tree-dump
+	directives.
+
 2015-07-12  Tom de Vries  <tom@codesourcery.com>
 
 	* c-c++-common/goacc/kernels-loop-nest.c: New test.
diff --git gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
index 3e06c9f..e8eebaa 100644
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c
@@ -37,6 +37,3 @@ main (void)
 /* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */
 
 /* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
-
-/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
-/* { dg-final { cleanup-tree-dump "optimized" } } */


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [gomp4, committed] Handle nested loops in kernels regions

[Jakub Jelinek]

On Mon, Jul 13, 2015 at 10:19:56AM +0200, Thomas Schwinge wrote:
> > We rely on pass_lim to move the *.omp_data_i loads out of the loop nest. 
> > For the test-case, pass_lim was managing to move the load out of the 
> > inner loop, but not the outer loop, because the load was classified as 
> > 'MOVE_PRESERVE_EXECUTION'. By marking the *.omp_data_i load 
> > non-trapping, it's now classified as 'MOVE_POSSIBLE', and moved out of 
> > the loop nest.
> 
> Should this go into trunk already?  (Jakub?)

I think so.

> Do we need to audit the
> code for constructs that need similar treatment?

That might be helpful.

	Jakub

[Committed] Mark *.omp_data_i as non-trapping

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 745 bytes --]

[ was: Re: [gomp4, committed] Handle nested loops in kernels regions ]

On 13/07/15 10:36, Jakub Jelinek wrote:
> On Mon, Jul 13, 2015 at 10:19:56AM +0200, Thomas Schwinge wrote:
>>> We rely on pass_lim to move the *.omp_data_i loads out of the loop nest.
>>> For the test-case, pass_lim was managing to move the load out of the
>>> inner loop, but not the outer loop, because the load was classified as
>>> 'MOVE_PRESERVE_EXECUTION'. By marking the *.omp_data_i load
>>> non-trapping, it's now classified as 'MOVE_POSSIBLE', and moved out of
>>> the loop nest.
>>
>> Should this go into trunk already?  (Jakub?)
>
> I think so.
>

I already bootstrapped and regtested (on x86_64) the patch on top of trunk.

Committed to trunk.

Thanks,
- Tom


[-- Attachment #2: 0001-Mark-.omp_data_i-as-non-trapping.patch --]
[-- Type: text/x-patch, Size: 610 bytes --]

Mark *.omp_data_i as non-trapping

2015-07-12  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (build_receiver_ref): Mark *.omp_data_i as non-trapping.
---
 gcc/omp-low.c | 1 +
 1 file changed, 1 insertion(+)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 22848a0..3135606 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1127,6 +1127,7 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
     field = x;
 
   x = build_simple_mem_ref (ctx->receiver_decl);
+  TREE_THIS_NOTRAP (x) = 1;
   x = omp_build_component_ref (x, field);
   if (by_ref)
     x = build_simple_mem_ref (x);
-- 
1.9.1

[PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 894 bytes --]

[ was: Re: [Committed] Mark *.omp_data_i as non-trapping ]

On 13/07/15 11:49, Tom de Vries wrote:
> [ was: Re: [gomp4, committed] Handle nested loops in kernels regions ]
>
> On 13/07/15 10:36, Jakub Jelinek wrote:
>> On Mon, Jul 13, 2015 at 10:19:56AM +0200, Thomas Schwinge wrote:
>>>> We rely on pass_lim to move the *.omp_data_i loads out of the loop
>>>> nest.
>>>> For the test-case, pass_lim was managing to move the load out of the
>>>> inner loop, but not the outer loop, because the load was classified as
>>>> 'MOVE_PRESERVE_EXECUTION'. By marking the *.omp_data_i load
>>>> non-trapping, it's now classified as 'MOVE_POSSIBLE', and moved out of
>>>> the loop nest.

This follow-up patch also marks the 'by_ref' mem_ref in 
build_receiver_ref as non-trapping.

Bootstrapped and reg-tested on x86_64.

OK for stage3 (because it's needed for the oacc kernels support)?

Thanks,
- Tom

[-- Attachment #2: 0001-Mark-by_ref-mem_ref-in-build_receiver_ref-as-non-trapping.patch --]
[-- Type: text/x-patch, Size: 681 bytes --]

Mark by_ref mem_ref in build_receiver_ref as non-trapping

2015-11-21  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (build_receiver_ref): Mark by_ref mem_ref as non-trapping.

---
 gcc/omp-low.c | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 830db75..78f2853 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1249,7 +1249,10 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
   TREE_THIS_NOTRAP (x) = 1;
   x = omp_build_component_ref (x, field);
   if (by_ref)
-    x = build_simple_mem_ref (x);
+    {
+      x = build_simple_mem_ref (x);
+      TREE_THIS_NOTRAP (x) = 1;
+    }
 
   return x;
 }

Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping

[Jakub Jelinek]

On Sat, Nov 21, 2015 at 07:34:02PM +0100, Tom de Vries wrote:
> Mark by_ref mem_ref in build_receiver_ref as non-trapping
> 
> 2015-11-21  Tom de Vries  <tom@codesourcery.com>
> 
> 	* omp-low.c (build_receiver_ref): Mark by_ref mem_ref as non-trapping.

This is ok.
> 
> ---
>  gcc/omp-low.c | 5 ++++-
>  1 file changed, 4 insertions(+), 1 deletion(-)
> 
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index 830db75..78f2853 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -1249,7 +1249,10 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
>    TREE_THIS_NOTRAP (x) = 1;
>    x = omp_build_component_ref (x, field);
>    if (by_ref)
> -    x = build_simple_mem_ref (x);
> +    {
> +      x = build_simple_mem_ref (x);
> +      TREE_THIS_NOTRAP (x) = 1;
> +    }
>  
>    return x;
>  }


	Jakub

Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping

[Richard Biener]

On Mon, Nov 23, 2015 at 9:45 AM, Jakub Jelinek <jakub@redhat.com> wrote:
> On Sat, Nov 21, 2015 at 07:34:02PM +0100, Tom de Vries wrote:
>> Mark by_ref mem_ref in build_receiver_ref as non-trapping
>>
>> 2015-11-21  Tom de Vries  <tom@codesourcery.com>
>>
>>       * omp-low.c (build_receiver_ref): Mark by_ref mem_ref as non-trapping.
>
> This is ok.

Are you sure this is properly re-set by inlining via

          /* We cannot propagate the TREE_THIS_NOTRAP flag if we have
             remapped a parameter as the property might be valid only
             for the parameter itself.  */
          if (TREE_THIS_NOTRAP (old)
              && (!is_parm (TREE_OPERAND (old, 0))
                  || (!id->transform_parameter && is_parm (ptr))))
            TREE_THIS_NOTRAP (*tp) = 1;

?  Or is this never hoistable to a place where TREE_THIS_NOTRAP is not true
even after inlining?  (I presume this is not directly a load via the
static chain pointer?)

>>
>> ---
>>  gcc/omp-low.c | 5 ++++-
>>  1 file changed, 4 insertions(+), 1 deletion(-)
>>
>> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
>> index 830db75..78f2853 100644
>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -1249,7 +1249,10 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
>>    TREE_THIS_NOTRAP (x) = 1;
>>    x = omp_build_component_ref (x, field);
>>    if (by_ref)
>> -    x = build_simple_mem_ref (x);
>> +    {
>> +      x = build_simple_mem_ref (x);
>> +      TREE_THIS_NOTRAP (x) = 1;
>> +    }
>>
>>    return x;
>>  }
>
>
>         Jakub

Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping

[Jakub Jelinek]

On Mon, Nov 23, 2015 at 11:39:26AM +0100, Richard Biener wrote:
> On Mon, Nov 23, 2015 at 9:45 AM, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Sat, Nov 21, 2015 at 07:34:02PM +0100, Tom de Vries wrote:
> >> Mark by_ref mem_ref in build_receiver_ref as non-trapping
> >>
> >> 2015-11-21  Tom de Vries  <tom@codesourcery.com>
> >>
> >>       * omp-low.c (build_receiver_ref): Mark by_ref mem_ref as non-trapping.
> >
> > This is ok.
> 
> Are you sure this is properly re-set by inlining via
> 
>           /* We cannot propagate the TREE_THIS_NOTRAP flag if we have
>              remapped a parameter as the property might be valid only
>              for the parameter itself.  */
>           if (TREE_THIS_NOTRAP (old)
>               && (!is_parm (TREE_OPERAND (old, 0))
>                   || (!id->transform_parameter && is_parm (ptr))))
>             TREE_THIS_NOTRAP (*tp) = 1;
> 
> ?  Or is this never hoistable to a place where TREE_THIS_NOTRAP is not true
> even after inlining?  (I presume this is not directly a load via the
> static chain pointer?)

I don't think inlining is ever around here, this is inside of the outlined
bodies of the OpenMP constructs, those are the *.omp_fn* artificial
functions called from libgomp, and is used in cases where
  .omp_data_i->field
is not the field itself, but pointer to the original variable.  The caller
of the libgomp functions that in the end invoke the .omp_fn* functions
guarantees that the field in that case is initialized to an address of the
original variables, is not NULL or some invalid pointer.

	Jakub