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. 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
