[gomp4] remap variables inside gang, worker, vector and collapse clauses

On 09/23/2015 10:42 AM, Cesar Philippidis wrote:

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

This patch, that is.

Cesar

Hi!

On Wed, 23 Sep 2015 10:57:40 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 09/23/2015 10:42 AM, Cesar Philippidis wrote:
> | Gang, worker, vector and collapse all contain optional arguments which
> | may be used during loop expansion. In OpenACC, those expressions could
> | contain variables

I'm fairly sure that at least the collapse clause needs to be a
compile-time constant?

> | but those variables aren't always getting remapped
> | automatically. This patch remaps those variables inside lower_omp_loop.

Shouldn't that be done in lower_rec_input_clauses?  (Maybe I'm confused
-- it's been a long time that I looked at this code.)  (Jakub?)

> | Note that I didn't need to use a tree walker for more complicated
> | expressions because it's not required. By the time those clauses reach
> | lower_omp_loop, only the result of the expression is available. So the
> | other variables in those expressions get remapped with everything else
> | during omplow. Therefore, the only problematic case is when the the
> | optional expression is just a decl, e.g. gang(static:foo).
> 
> > I've applied this patch to gomp-4_0-branch.

> 2015-09-23  Cesar Philippidis  <cesar@codesourcery.com>
> 
> 	gcc/
> 	* omp-low.c (lower_omp_for): Remap any variables present in
> 	OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR and
> 	OMP_CLAUSE_COLLAPSE becuase they will be used later by expand_omp_for.
> 
> 	libgomp/
> 	* testsuite/libgomp.oacc-c-c++-common/gang-static-2.c: Test if
> 	static gang expressions containing variables work.
> 	* testsuite/libgomp.oacc-fortran/gang-static-1.f90: Likewise.
> 
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index ec76096..3f36b7a 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -11325,6 +11325,35 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
>    if (oacc_tail)
>      gimple_seq_add_seq (&body, oacc_tail);
>  
> +  /* Update the variables inside any clauses which may be involved in loop
> +     expansion later on.  */
> +  for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
> +    {
> +      int args;
> +
> +      switch (OMP_CLAUSE_CODE (c))
> +	{
> +	default:
> +	  args = 0;
> +	  break;
> +	case OMP_CLAUSE_GANG:
> +	  args = 2;
> +	  break;
> +	case OMP_CLAUSE_VECTOR:
> +	case OMP_CLAUSE_WORKER:
> +	case OMP_CLAUSE_COLLAPSE:
> +	  args = 1;
> +	  break;
> +	}
> +
> +      for (int i = 0; i < args; i++)
> +	{
> +	  tree expr = OMP_CLAUSE_OPERAND (c, i);
> +	  if (expr && DECL_P (expr))
> +	    OMP_CLAUSE_OPERAND (c, i) = build_outer_var_ref (expr, ctx);
> +	}
> +    }
> +
>    pop_gimplify_context (new_stmt);
>  
>    gimple_bind_append_vars (new_stmt, ctx->block_vars);
> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
> index 3a9a508..20a866d 100644
> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
> @@ -39,7 +39,7 @@ int
>  main ()
>  {
>    int a[N];
> -  int i;
> +  int i, x;
>  
>  #pragma acc parallel loop gang (static:*) num_gangs (10)
>    for (i = 0; i < 100; i++)
> @@ -78,5 +78,21 @@ main ()
>  
>    test_nonstatic (a, 10);
>  
> +  /* Static arguments with a variable expression.  */
> +
> +  x = 20;
> +#pragma acc parallel loop gang (static:0+x) num_gangs (10)
> +  for (i = 0; i < 100; i++)
> +    a[i] = GANG_ID (i);
> +
> +  test_static (a, 10, 20);
> +
> +  x = 20;
> +#pragma acc parallel loop gang (static:x) num_gangs (10)
> +  for (i = 0; i < 100; i++)
> +    a[i] = GANG_ID (i);
> +
> +  test_static (a, 10, 20);
> +
>    return 0;
>  }
> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/gang-static-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/gang-static-1.f90
> index e562535..7d56060 100644
> --- a/libgomp/testsuite/libgomp.oacc-fortran/gang-static-1.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/gang-static-1.f90
> @@ -3,6 +3,7 @@
>  program main
>    integer, parameter :: n = 100
>    integer i, a(n), b(n)
> +  integer x
>  
>    do i = 1, n
>       b(i) = i
> @@ -48,6 +49,23 @@ program main
>  
>    call test (a, b, 20, n)
>  
> +  x = 5
> +  !$acc parallel loop gang (static:0+x) num_gangs (10)
> +  do i = 1, n
> +     a(i) = b(i) + 5
> +  end do
> +  !$acc end parallel loop
> +
> +  call test (a, b, 5, n)
> +
> +  x = 10
> +  !$acc parallel loop gang (static:x) num_gangs (10)
> +  do i = 1, n
> +     a(i) = b(i) + 10
> +  end do
> +  !$acc end parallel loop
> +
> +  call test (a, b, 10, n)
>  end program main
>  
>  subroutine test (a, b, sarg, n)


Grüße,
 Thomas