static splay_tree all_contexts;
static int taskreg_nesting_level;
static int target_nesting_level;
-static bitmap task_shared_vars;
+static bitmap make_addressable_vars;
static bitmap global_nonaddressable_vars;
static vec<omp_context *> taskreg_contexts;
static vec<gomp_task *> task_cpyfns;
/* Taking address of OUTER in lower_send_shared_vars
might need regimplification of everything that uses the
variable. */
- if (!task_shared_vars)
- task_shared_vars = BITMAP_ALLOC (NULL);
- bitmap_set_bit (task_shared_vars, DECL_UID (outer));
+ if (!make_addressable_vars)
+ make_addressable_vars = BITMAP_ALLOC (NULL);
+ bitmap_set_bit (make_addressable_vars, DECL_UID (outer));
TREE_ADDRESSABLE (outer) = 1;
}
return true;
else
record_vars (copy);
- /* If VAR is listed in task_shared_vars, it means it wasn't
- originally addressable and is just because task needs to take
- it's address. But we don't need to take address of privatizations
+ /* If VAR is listed in make_addressable_vars, it wasn't
+ originally addressable, but was only later made so.
+ We don't need to take address of privatizations
from that var. */
if (TREE_ADDRESSABLE (var)
- && ((task_shared_vars
- && bitmap_bit_p (task_shared_vars, DECL_UID (var)))
+ && ((make_addressable_vars
+ && bitmap_bit_p (make_addressable_vars, DECL_UID (var)))
|| (global_nonaddressable_vars
&& bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var)))))
TREE_ADDRESSABLE (copy) = 0;
gcc_checking_assert (DECL_P (decl));
gcc_checking_assert (!TREE_ADDRESSABLE (decl));
+ if (!make_addressable_vars)
+ make_addressable_vars = BITMAP_ALLOC (NULL);
+ bitmap_set_bit (make_addressable_vars, DECL_UID (decl));
TREE_ADDRESSABLE (decl) = 1;
if (dump_enabled_p ())
if (ctx->record_type == NULL_TREE)
return;
- /* If any task_shared_vars were needed, verify all
+ /* If any make_addressable_vars were needed, verify all
OMP_CLAUSE_SHARED clauses on GIMPLE_OMP_{PARALLEL,TASK,TEAMS}
statements if use_pointer_for_field hasn't changed
because of that. If it did, update field types now. */
- if (task_shared_vars)
+ if (make_addressable_vars)
{
tree c;
the receiver side will use them directly. */
if (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
continue;
- if (!bitmap_bit_p (task_shared_vars, DECL_UID (decl))
+ if (!bitmap_bit_p (make_addressable_vars, DECL_UID (decl))
|| !use_pointer_for_field (decl, ctx))
continue;
tree field = lookup_field (decl, ctx);
/* Callback for lower_omp_1. Return non-NULL if *tp needs to be
regimplified. If DATA is non-NULL, lower_omp_1 is outside
- of OMP context, but with task_shared_vars set. */
+ of OMP context, but with make_addressable_vars set. */
static tree
lower_omp_regimplify_p (tree *tp, int *walk_subtrees,
&& DECL_HAS_VALUE_EXPR_P (t))
return t;
- if (task_shared_vars
+ if (make_addressable_vars
&& DECL_P (t)
- && bitmap_bit_p (task_shared_vars, DECL_UID (t)))
+ && bitmap_bit_p (make_addressable_vars, DECL_UID (t)))
return t;
/* If a global variable has been privatized, TREE_CONSTANT on
if (gimple_has_location (stmt))
input_location = gimple_location (stmt);
- if (task_shared_vars)
+ if (make_addressable_vars)
memset (&wi, '\0', sizeof (wi));
/* If we have issued syntax errors, avoid doing any heavy lifting.
case GIMPLE_COND:
{
gcond *cond_stmt = as_a <gcond *> (stmt);
- if ((ctx || task_shared_vars)
+ if ((ctx || make_addressable_vars)
&& (walk_tree (gimple_cond_lhs_ptr (cond_stmt),
lower_omp_regimplify_p,
ctx ? NULL : &wi, NULL)
lower_omp_critical (gsi_p, ctx);
break;
case GIMPLE_OMP_ATOMIC_LOAD:
- if ((ctx || task_shared_vars)
+ if ((ctx || make_addressable_vars)
&& walk_tree (gimple_omp_atomic_load_rhs_ptr (
as_a <gomp_atomic_load *> (stmt)),
lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
default:
regimplify:
- if ((ctx || task_shared_vars)
+ if ((ctx || make_addressable_vars)
&& walk_gimple_op (stmt, lower_omp_regimplify_p,
ctx ? NULL : &wi))
{
if (all_contexts->root)
{
- if (task_shared_vars)
+ if (make_addressable_vars)
push_gimplify_context ();
lower_omp (&body, NULL);
- if (task_shared_vars)
+ if (make_addressable_vars)
pop_gimplify_context (NULL);
}
splay_tree_delete (all_contexts);
all_contexts = NULL;
}
- BITMAP_FREE (task_shared_vars);
+ BITMAP_FREE (make_addressable_vars);
BITMAP_FREE (global_nonaddressable_vars);
/* If current function is a method, remove artificial dummy VAR_DECL created
/* { dg-additional-options "--param openacc-kernels=decompose" } */
-/* { dg-additional-options "-fchecking" }
- { dg-ice TODO }
- { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
- { dg-prune-output {during GIMPLE pass: lower} } */
-
/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
- { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail *-*-* } 0 }
{ dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */
/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k1 } */
/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
- /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
- /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {w/o debug} { target *-*-* } l_loop_k1 }
+ { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
for (k = 0; k < 2; k++)
arr_0 += k;
- /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
}
}
/* { dg-additional-options "--param openacc-kernels=decompose" } */
-/* { dg-additional-options "-fchecking" }
- { dg-ice TODO }
- { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
- { dg-prune-output {during GIMPLE pass: lower} } */
-
-/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first.
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */
/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
/* { dg-additional-options "-fopt-info-all-omp" } */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k1 } */
- /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
- /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
- /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
- /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } l_loop_k1 } */
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {w/o debug} { target *-*-* } l_loop_k1 }
+ { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
for (k = 0; k < 2; k++)
arr_0 += k;
- /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
}
}
/* { dg-additional-options "--param openacc-kernels=decompose" } */
-/* { dg-additional-options "-fchecking" }
- { dg-ice TODO }
- { dg-prune-output {k = 0 \+ \.offset\.[0-9]+;} }
- { dg-prune-output {k = 0 \+ 2;} }
- { dg-prune-output {during IPA pass: \*free_lang_data} } */
-
/* { dg-additional-options "-fopt-info-all-omp" } */
/* { dg-additional-options "--param=openacc-privatization=noisy" }
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k1 } */
/* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
for (k = 0; k < 2; k++)
arr_0 = k;
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k2 } */
/* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k2 } */
for (k = 0; k < 2; k++)
arr_0 = k;
}
}
-/* { dg-bogus {error: non-register as LHS of binary operation} {} { xfail *-*-* } .-1 } */
/* { dg-additional-options "--param openacc-kernels=decompose" } */
-/* { dg-additional-options "-fchecking" }
- { dg-ice TODO }
- { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
- { dg-prune-output {D\.[0-9]+ = arr_0\.1 \+ k;} }
- { dg-prune-output {during GIMPLE pass: lower} } */
-
/* { dg-additional-options "-fopt-info-all-omp" } */
/* { dg-additional-options "--param=openacc-privatization=noisy" }
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k1 } */
/* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k1 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
for (k = 0; k < 2; k++)
arr_0 += k;
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
#pragma acc loop /* { dg-line l_loop_k2 } */
/* { dg-note {variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_k2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k2 } */
for (k = 0; k < 2; k++)
arr_0 += k;
- /* { dg-bogus {error: invalid operands in binary operation} {} { xfail *-*-* } .-1 } */
}
}
/* { dg-additional-options "--param=openacc-privatization=noisy" }
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
- for testing/documenting aspects of that functionality. */
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
passed to 'incr' may be unset, and in that case, it will be set to [...]",
so to maintain compatibility with earlier Tcl releases, we manually
initialize counter variables:
- { dg-line l_dummy[variable c_compute 0 c_loop_i 0] }
+ { dg-line l_dummy[variable c_compute 0 c_loop_c 0 c_loop_i 0] }
{ dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
"WARNING: dg-line var l_dummy defined, but not used". */
#undef NDEBUG
#include <assert.h>
+static int g1;
+static int g2;
+
int main()
{
int a = 0;
(volatile int *) &a;
#define N 123
int b[N] = { 0 };
+ unsigned long long f1;
+ /*TODO See above. */
+ (volatile void *) &f1;
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'g2\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
{
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
int c = 234;
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
a = c;
+
+ /* PR104132, PR104133 */
+ {
+ /* Use the 'kernels'-top-level 'int c' as loop variable. */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
+ for (c = 0; c < N / 2; c++)
+ b[c] -= 10;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
+ for (c = 0; c < N / 2; c++)
+ g1 = c;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
+ for (c = 0; c <= N; c++)
+ g2 += c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ f1 = 1;
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c$c_loop_c } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
+ for (c = 20; c > 0; --c)
+ f1 *= c;
+
+ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ if (c != 234)
+ __builtin_abort ();
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
+ }
}
- for (int i = 0; i < N; ++i)
- assert (b[i] == 234);
assert (a == 234);
+ for (int i = 0; i < N; ++i)
+ if (i < N / 2)
+ assert (b[i] == 234 - 10);
+ else
+ assert (b[i] == 234);
+ assert (g1 == N / 2 - 1);
+ assert (g2 == N * (N + 1) / 2);
+ assert (f1 == 2432902008176640000ULL);
return 0;
}