On 13-06-19 10:34, Jakub Jelinek wrote:
> Hi!
>
> The OpenMP specification isn't clear on this, I'll work on getting that
> clarified for 5.1, but the agreement on omp-lang has been that it should
> work the way the patch implements it, static block scope variables inside of
> #pragma omp target or #pragma omp declare target routines are handled as if
> they have #pragma omp declare target to (variable).
>
> Bootstrapped/regtested on x86_64-linux and i686-linux, unfortunately it
> regresses:
> +FAIL: c-c++-common/goacc/routine-5.c (test for errors, line 204)
>
> Thus, I'm not committing it right now and want to ask what should be done
> for OpenACC.
OpenACC 2.6 - 2.15.1. Routine Directive - Restrictions:
...
In C and C++, function static variables are not supported in functions
to which a routine directive applies.
...
[ And text is still the same for 2.7. ]
Thanks,
- Tom
> The patch uses & ORT_TARGET tests, so it affects both OpenMP
> target region, and OpenACC parallel/kernels and both OpenMP and OpenACC
> target routines. Is it ok to do it that way and just adjust the routine-5.c
> test, or shall it test (ctx->region_type & (ORT_TARGET | ORT_ACC)) ==
> ORT_TARGET, i.e. only OpenMP and not OpenACC? If so, there is still the
> problem that gimplify_body.c does:
> if (flag_openacc || flag_openmp)
> {
> gcc_assert (gimplify_omp_ctxp == NULL);
> if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
> gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
> }
> We'd need different attribute (or additional attribute) for OpenACC routines
> and would need to use new_omp_context (cond ? ORT_TARGET : ORT_ACC_PARALLEL)
> or similar to express OpenACC routines.
>
> 2019-06-12 Jakub Jelinek <ja...@redhat.com>
>
> PR middle-end/90779
> * gimplify.c (gimplify_bind_expr): Add "omp declare target" attributes
> to static block scope variables inside of target region or target
> functions.
>
> * testsuite/libgomp.c/pr90779.c: New test.
> * testsuite/libgomp.fortran/pr90779.f90: New test.
>
> --- gcc/gimplify.c.jj 2019-06-10 19:42:03.868959986 +0200
> +++ gcc/gimplify.c 2019-06-12 13:00:18.765167777 +0200
> @@ -1323,17 +1323,37 @@ gimplify_bind_expr (tree *expr_p, gimple
> struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
>
> /* Mark variable as local. */
> - if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t)
> - && (! DECL_SEEN_IN_BIND_EXPR_P (t)
> - || splay_tree_lookup (ctx->variables,
> - (splay_tree_key) t) == NULL))
> + if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t))
> {
> - if (ctx->region_type == ORT_SIMD
> - && TREE_ADDRESSABLE (t)
> - && !TREE_STATIC (t))
> - omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
> - else
> - omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
> + if (! DECL_SEEN_IN_BIND_EXPR_P (t)
> + || splay_tree_lookup (ctx->variables,
> + (splay_tree_key) t) == NULL)
> + {
> + if (ctx->region_type == ORT_SIMD
> + && TREE_ADDRESSABLE (t)
> + && !TREE_STATIC (t))
> + omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN);
> + else
> + omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN);
> + }
> + /* Static locals inside of target construct or offloaded
> + routines need to be "omp declare target". */
> + if (TREE_STATIC (t))
> + for (; ctx; ctx = ctx->outer_context)
> + if ((ctx->region_type & ORT_TARGET) != 0)
> + {
> + if (!lookup_attribute ("omp declare target",
> + DECL_ATTRIBUTES (t)))
> + {
> + tree id = get_identifier ("omp declare target");
> + DECL_ATTRIBUTES (t)
> + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
> + varpool_node *node = varpool_node::get (t);
> + if (node)
> + node->offloadable = 1;
> + }
> + break;
> + }
> }
>
> DECL_SEEN_IN_BIND_EXPR_P (t) = 1;
> --- libgomp/testsuite/libgomp.c/pr90779.c.jj 2019-06-12 13:01:57.081667587
> +0200
> +++ libgomp/testsuite/libgomp.c/pr90779.c 2019-06-12 12:41:15.637730797
> +0200
> @@ -0,0 +1,18 @@
> +/* PR middle-end/90779 */
> +
> +extern void abort (void);
> +
> +int
> +main ()
> +{
> + int i, j;
> + for (i = 0; i < 2; ++i)
> + #pragma omp target map(from: j)
> + {
> + static int k = 5;
> + j = ++k;
> + }
> + if (j != 7)
> + abort ();
> + return 0;
> +}
> --- libgomp/testsuite/libgomp.fortran/pr90779.f90.jj 2019-06-12
> 12:43:17.891825811 +0200
> +++ libgomp/testsuite/libgomp.fortran/pr90779.f90 2019-06-12
> 12:43:08.421973375 +0200
> @@ -0,0 +1,12 @@
> +! PR middle-end/90779
> +
> +program pr90779
> + implicit none
> + integer :: v(4), i
> +
> + !$omp target map(from:v)
> + v(:) = (/ (i, i=1,4) /)
> + !$omp end target
> +
> + if (any (v .ne. (/ (i, i=1,4) /))) stop 1
> +end program
>
> Jakub
>
