* [gomp4, committed] Handle nested loops in kernels regions
@ 2015-07-12 12:46 Tom de Vries
2015-07-13 8:20 ` Thomas Schwinge
0 siblings, 1 reply; 8+ messages in thread
From: Tom de Vries @ 2015-07-12 12:46 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
[-- 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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4, committed] Handle nested loops in kernels regions
2015-07-12 12:46 [gomp4, committed] Handle nested loops in kernels regions Tom de Vries
@ 2015-07-13 8:20 ` Thomas Schwinge
2015-07-13 8:36 ` Jakub Jelinek
0 siblings, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2015-07-13 8:20 UTC (permalink / raw)
To: Tom de Vries, gcc-patches, Jakub Jelinek
[-- 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 --]
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4, committed] Handle nested loops in kernels regions
2015-07-13 8:20 ` Thomas Schwinge
@ 2015-07-13 8:36 ` Jakub Jelinek
2015-07-13 9:49 ` [Committed] Mark *.omp_data_i as non-trapping Tom de Vries
0 siblings, 1 reply; 8+ messages in thread
From: Jakub Jelinek @ 2015-07-13 8:36 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Tom de Vries, gcc-patches
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Committed] Mark *.omp_data_i as non-trapping
2015-07-13 8:36 ` Jakub Jelinek
@ 2015-07-13 9:49 ` Tom de Vries
2015-11-21 18:49 ` [PATCH] Mark by_ref mem_ref in build_receiver_ref " Tom de Vries
0 siblings, 1 reply; 8+ messages in thread
From: Tom de Vries @ 2015-07-13 9:49 UTC (permalink / raw)
To: Jakub Jelinek, Thomas Schwinge; +Cc: gcc-patches
[-- 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
^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping
2015-07-13 9:49 ` [Committed] Mark *.omp_data_i as non-trapping Tom de Vries
@ 2015-11-21 18:49 ` Tom de Vries
2015-11-23 9:10 ` Jakub Jelinek
0 siblings, 1 reply; 8+ messages in thread
From: Tom de Vries @ 2015-11-21 18:49 UTC (permalink / raw)
To: Jakub Jelinek, Thomas Schwinge; +Cc: gcc-patches
[-- 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;
}
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping
2015-11-21 18:49 ` [PATCH] Mark by_ref mem_ref in build_receiver_ref " Tom de Vries
@ 2015-11-23 9:10 ` Jakub Jelinek
2015-11-23 10:40 ` Richard Biener
0 siblings, 1 reply; 8+ messages in thread
From: Jakub Jelinek @ 2015-11-23 9:10 UTC (permalink / raw)
To: Tom de Vries; +Cc: Thomas Schwinge, gcc-patches
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping
2015-11-23 9:10 ` Jakub Jelinek
@ 2015-11-23 10:40 ` Richard Biener
2015-11-23 11:02 ` Jakub Jelinek
0 siblings, 1 reply; 8+ messages in thread
From: Richard Biener @ 2015-11-23 10:40 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Tom de Vries, Thomas Schwinge, gcc-patches
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] Mark by_ref mem_ref in build_receiver_ref as non-trapping
2015-11-23 10:40 ` Richard Biener
@ 2015-11-23 11:02 ` Jakub Jelinek
0 siblings, 0 replies; 8+ messages in thread
From: Jakub Jelinek @ 2015-11-23 11:02 UTC (permalink / raw)
To: Richard Biener; +Cc: Tom de Vries, Thomas Schwinge, gcc-patches
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
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2015-11-23 10:55 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-07-12 12:46 [gomp4, committed] Handle nested loops in kernels regions Tom de Vries
2015-07-13 8:20 ` Thomas Schwinge
2015-07-13 8:36 ` Jakub Jelinek
2015-07-13 9:49 ` [Committed] Mark *.omp_data_i as non-trapping Tom de Vries
2015-11-21 18:49 ` [PATCH] Mark by_ref mem_ref in build_receiver_ref " Tom de Vries
2015-11-23 9:10 ` Jakub Jelinek
2015-11-23 10:40 ` Richard Biener
2015-11-23 11:02 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).