diff mbox

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

Message ID 87ioulgisw.fsf@schwinge.name
State New
Headers show

Commit Message

Thomas Schwinge Dec. 19, 2013, 11:51 a.m. UTC 
Hi!

On Wed, 18 Dec 2013 22:46:48 +0100, Jakub Jelinek <jakub@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 <jakub@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 <thomas@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.



Grüße,
 Thomas

Comments

Jakub Jelinek Dec. 19, 2013, 12:16 p.m. UTC | #1 
On Thu, Dec 19, 2013 at 12:51:43PM +0100, Thomas Schwinge wrote:
> OK for gomp-4_0-branch, and trunk (without the first hunk, obviously)?
> 
> commit 1f4dbb1842804b39c3b7ac1e80783734516dc965
> Author: Thomas Schwinge <thomas@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.

Sure.
> --- 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);
> 
> 

	Jakub
diff mbox

Patch

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);