Re: GOMP_target: alignment (was: [gomp4] #pragma omp target* fixes)

[Thomas Schwinge]

Hi!

On Wed, 18 Dec 2013 22:46:48 +0100, Jakub Jelinek <ja...@redhat.com> wrote:
> On Wed, Dec 18, 2013 at 09:03:40PM +0100, Thomas Schwinge wrote:
> > On Mon, 16 Dec 2013 16:38:18 +0100, Jakub Jelinek <ja...@redhat.com> wrote:
> > > The reason for 3 separate arrays is that some of the values
> > > are always variable, some are sometimes variable (sizes), some are
> > > never variable (alignment + kind).
> > 
> > Related to this, in gcc/omp-low.c:lower_omp_target, I see:
> > 
> >           tree clobber = build_constructor (ctx->record_type, NULL);
> >           TREE_THIS_VOLATILE (clobber) = 1;
> >           gimple_seq_add_stmt (&olist, gimple_build_assign 
> > (ctx->sender_decl,
> >                                                             clobber));
> 
> Clobber stmt is [...]

Thanks for explaining, and basically confirming my assumption.  :-)

On Wed, 18 Dec 2013 21:03:40 +0100, I wrote:
> And, why doesn't the same also need to be done for the sizes object (in
> the non-static case)?

OK for gomp-4_0-branch, and trunk (without the first hunk, obviously)?

commit 1f4dbb1842804b39c3b7ac1e80783734516dc965
Author: Thomas Schwinge <tho...@codesourcery.com>
Date:   Thu Dec 19 12:41:21 2013 +0100

    Add clobber for object, after last use.
    
        gcc/
        * omp-low.c (lower_oacc_parallel, lower_omp_target): Add clobber
        for sizes array, after last use.

diff --git gcc/omp-low.c gcc/omp-low.c
index d3e22a1..f17affe 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -8917,6 +8917,13 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, 
omp_context *ctx)
                                        TREE_VEC_ELT (t, 1)),
                                &initlist, true, NULL_TREE);
          gimple_seq_add_seq (&ilist, initlist);
+
+         tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)),
+                                           NULL);
+         TREE_THIS_VOLATILE (clobber) = 1;
+         gimple_seq_add_stmt (&olist,
+                              gimple_build_assign (TREE_VEC_ELT (t, 1),
+                                                   clobber));
        }
 
       tree clobber = build_constructor (ctx->record_type, NULL);
@@ -10412,6 +10419,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, 
omp_context *ctx)
                                        TREE_VEC_ELT (t, 1)),
                                &initlist, true, NULL_TREE);
          gimple_seq_add_seq (&ilist, initlist);
+
+         tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)),
+                                           NULL);
+         TREE_THIS_VOLATILE (clobber) = 1;
+         gimple_seq_add_stmt (&olist,
+                              gimple_build_assign (TREE_VEC_ELT (t, 1),
+                                                   clobber));
        }
 
       tree clobber = build_constructor (ctx->record_type, NULL);


Grüße,
 Thomas

pgpw8eRmXmwmb.pgp
Description: PGP signature