OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133]
... by generalizing the existing 'gcc/omp-low.cc:task_shared_vars'.
Fix-up for commit 9b32c1669a
"OpenACC 'kernels' decomposition: Mark variables used in
synthesized data clauses as addressable [PR100280]".
PR middle-end/100280
PR middle-end/104132
PR middle-end/104133
gcc/
* omp-low.cc (task_shared_vars): Rename to
'make_addressable_vars'. Adjust all users.
(scan_sharing_clauses) <OMP_CLAUSE_MAP> Use it for
'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs, too.
gcc/testsuite/
* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Adjust.
* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
Extend.
This commit is contained in:
Thomas Schwinge
parent
de6e81ea96
commit
8935589b49
6 changed files with 95 additions and 60 deletions
|
|
@ -188,7 +188,7 @@ struct omp_context
|
|||
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;
|
||||
|
|
@ -572,9 +572,9 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
|
|||
/* 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;
|
||||
|
|
@ -601,13 +601,13 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx)
|
|||
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;
|
||||
|
|
@ -1502,6 +1502,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
|
|||
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 ())
|
||||
|
|
@ -2402,11 +2405,11 @@ finish_taskreg_scan (omp_context *ctx)
|
|||
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;
|
||||
|
||||
|
|
@ -2421,7 +2424,7 @@ finish_taskreg_scan (omp_context *ctx)
|
|||
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);
|
||||
|
|
@ -14071,7 +14074,7 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *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,
|
||||
|
|
@ -14085,9 +14088,9 @@ 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
|
||||
|
|
@ -14172,7 +14175,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|
|||
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.
|
||||
|
|
@ -14189,7 +14192,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|
|||
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)
|
||||
|
|
@ -14281,7 +14284,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|
|||
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))
|
||||
|
|
@ -14402,7 +14405,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|
|||
|
||||
default:
|
||||
regimplify:
|
||||
if ((ctx || task_shared_vars)
|
||||
if ((ctx || make_addressable_vars)
|
||||
&& walk_gimple_op (stmt, lower_omp_regimplify_p,
|
||||
ctx ? NULL : &wi))
|
||||
{
|
||||
|
|
@ -14466,10 +14469,10 @@ execute_lower_omp (void)
|
|||
|
||||
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);
|
||||
}
|
||||
|
||||
|
|
@ -14478,7 +14481,7 @@ execute_lower_omp (void)
|
|||
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
|
||||
|
|
|
|||
|
|
@ -1,12 +1,6 @@
|
|||
/* { 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. */
|
||||
|
||||
|
|
@ -35,11 +29,10 @@ foo (void)
|
|||
/* { 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 } */
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,11 +1,7 @@
|
|||
/* { 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" } */
|
||||
|
|
@ -32,12 +28,11 @@ foo (void)
|
|||
|
||||
/* { 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 } */
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,11 +1,5 @@
|
|||
/* { 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" }
|
||||
|
|
@ -27,14 +21,15 @@ foo (void)
|
|||
/* { 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 } */
|
||||
|
|
|
|||
|
|
@ -1,11 +1,5 @@
|
|||
/* { 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" }
|
||||
|
|
@ -29,14 +23,15 @@ foo (void)
|
|||
/* { 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 } */
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -7,19 +7,23 @@
|
|||
|
||||
/* { 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;
|
||||
|
|
@ -27,8 +31,12 @@ int main()
|
|||
(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;
|
||||
|
|
@ -46,11 +54,57 @@ int main()
|
|||
|
||||
/* { 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;
|
||||
}
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue