Hi! On Wed, 23 Sep 2015 10:57:40 -0700, Cesar Philippidis <ce...@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 <ce...@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
signature.asc
Description: PGP signature
