OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086]
... like in recent commit 9b32c1669a
"OpenACC 'kernels' decomposition: Mark variables used in synthesized
data clauses as addressable [PR100280]". Otherwise, we may run into
'gcc/omp-low.cc:lower_omp_target':
13125 else if (is_gimple_reg (var))
13126 {
13127 gcc_assert (offloaded);
PR middle-end/100280
PR middle-end/104086
gcc/
* omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1):
Mark variables used in 'present' clauses as addressable.
* omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Gracefully
handle duplicate 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'.
gcc/testsuite/
* c-c++-common/goacc/kernels-decompose-pr104086-1.c: Adjust,
extend.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c:
Merge this...
* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c:
..., and this...
* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: ... into
this, and adjust.
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
Extend.
This commit is contained in:
Thomas Schwinge
parent
9781ae3a25
commit
337ed336d7
7 changed files with 168 additions and 83 deletions
|
|
@ -1501,11 +1501,14 @@ 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;
|
||||
bool decl_addressable = TREE_ADDRESSABLE (decl);
|
||||
if (!decl_addressable)
|
||||
{
|
||||
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 ())
|
||||
{
|
||||
|
|
@ -1517,10 +1520,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
|
|||
# pragma GCC diagnostic push
|
||||
# pragma GCC diagnostic ignored "-Wformat"
|
||||
#endif
|
||||
dump_printf_loc (MSG_NOTE, d_u_loc,
|
||||
"variable %<%T%>"
|
||||
" made addressable\n",
|
||||
decl);
|
||||
if (!decl_addressable)
|
||||
dump_printf_loc (MSG_NOTE, d_u_loc,
|
||||
"variable %<%T%>"
|
||||
" made addressable\n",
|
||||
decl);
|
||||
else
|
||||
dump_printf_loc (MSG_NOTE, d_u_loc,
|
||||
"variable %<%T%>"
|
||||
" already made addressable\n",
|
||||
decl);
|
||||
#if __GNUC__ >= 10
|
||||
# pragma GCC diagnostic pop
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -1468,6 +1468,38 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
|
|||
/* Now that this data is mapped, turn the data clause on the
|
||||
inner OpenACC 'kernels' into a 'present' clause. */
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
|
||||
|
||||
/* See <https://gcc.gnu.org/PR100280>,
|
||||
<https://gcc.gnu.org/PR104086>. */
|
||||
if (DECL_P (decl)
|
||||
&& !TREE_ADDRESSABLE (decl))
|
||||
{
|
||||
/* Request that OMP lowering make 'decl' addressable. */
|
||||
OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1;
|
||||
|
||||
if (dump_enabled_p ())
|
||||
{
|
||||
location_t loc = OMP_CLAUSE_LOCATION (new_clause);
|
||||
const dump_user_location_t d_u_loc
|
||||
= dump_user_location_t::from_location_t (loc);
|
||||
/* PR100695 "Format decoder, quoting in 'dump_printf'
|
||||
etc." */
|
||||
#if __GNUC__ >= 10
|
||||
# pragma GCC diagnostic push
|
||||
# pragma GCC diagnostic ignored "-Wformat"
|
||||
#endif
|
||||
dump_printf_loc
|
||||
(MSG_NOTE, d_u_loc,
|
||||
"OpenACC %<kernels%> decomposition:"
|
||||
" variable %<%T%> in %qs clause"
|
||||
" requested to be made addressable\n",
|
||||
decl,
|
||||
user_omp_clause_code_name (new_clause, true));
|
||||
#if __GNUC__ >= 10
|
||||
# pragma GCC diagnostic pop
|
||||
#endif
|
||||
}
|
||||
}
|
||||
}
|
||||
break;
|
||||
|
||||
|
|
|
|||
|
|
@ -1,8 +1,5 @@
|
|||
/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c'. */
|
||||
|
||||
/* { dg-additional-options "-fchecking" }
|
||||
{ dg-ice TODO }
|
||||
{ dg-prune-output {during GIMPLE pass: omplower} } */
|
||||
/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c', and then
|
||||
extended. */
|
||||
|
||||
/* { dg-additional-options "--param openacc-kernels=decompose" } */
|
||||
|
||||
|
|
@ -14,12 +11,38 @@ void
|
|||
foo (void)
|
||||
{
|
||||
#pragma acc data /* { dg-line l_data1 } */
|
||||
/* { dg-bogus {note: variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO 'data'} { xfail *-*-* } l_data1 } */
|
||||
/* { dg-bogus {note: variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } l_data1 } */
|
||||
{
|
||||
int i;
|
||||
|
||||
#pragma acc kernels
|
||||
#pragma acc kernels /* { dg-line l_compute1 } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 }
|
||||
{ dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i = 0;
|
||||
|
||||
#pragma acc kernels /* { dg-line l_compute2 } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute2 }
|
||||
{ dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute2 } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i = -1;
|
||||
}
|
||||
}
|
||||
|
||||
void
|
||||
foo2 (void)
|
||||
{
|
||||
int i[1];
|
||||
|
||||
#pragma acc kernels /* { dg-line l2_compute1 } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute1 }
|
||||
{ dg-note {variable 'i' made addressable} {} { target *-*-* } l2_compute1 } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i[0] = 0;
|
||||
|
||||
#pragma acc kernels /* { dg-line l2_compute2 } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute2 }
|
||||
{ dg-note {variable 'i' already made addressable} {} { target *-*-* } l2_compute2 } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i[0] = -1;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,22 +0,0 @@
|
|||
/* { dg-additional-options "--param=openacc-kernels=decompose" } */
|
||||
/* ICE similar to PR100280, but not the same.
|
||||
{ dg-ice "TODO" }
|
||||
TODO { dg-prune-output "during GIMPLE pass: omplower" }
|
||||
TODO { dg-do link } */
|
||||
|
||||
/* { dg-additional-options "-fopt-info-omp-all" }
|
||||
{ dg-additional-options "-foffload=-fopt-info-all-omp" } */
|
||||
|
||||
/* { dg-additional-options "--param=openacc-privatization=noisy" }
|
||||
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
|
||||
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} } */
|
||||
|
||||
#undef KERNELS_DECOMPOSE_ICE_HACK
|
||||
#include "declare-vla.c"
|
||||
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */
|
||||
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */
|
||||
|
||||
/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */
|
||||
|
|
@ -1,29 +0,0 @@
|
|||
/* { dg-additional-options "--param=openacc-kernels=decompose" } */
|
||||
|
||||
/* See also 'declare-vla-kernels-decompose-ice-1.c'. */
|
||||
|
||||
/* { dg-additional-options "-fopt-info-omp-all" }
|
||||
{ dg-additional-options "-foffload=-fopt-info-all-omp" } */
|
||||
|
||||
/* { dg-additional-options "--param=openacc-privatization=noisy" }
|
||||
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
|
||||
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} } */
|
||||
|
||||
#define KERNELS_DECOMPOSE_ICE_HACK
|
||||
#include "declare-vla.c"
|
||||
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */
|
||||
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */
|
||||
|
||||
/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */
|
||||
|
||||
/* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 }
|
||||
{ dg-note {variable 'N\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 } */
|
||||
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 24 }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 24 } */
|
||||
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 58 }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 58 } */
|
||||
|
|
@ -1,5 +1,7 @@
|
|||
/* Verify OpenACC 'declare' with VLAs. */
|
||||
|
||||
/* { dg-additional-options "--param=openacc-kernels=decompose" } */
|
||||
|
||||
/* { dg-additional-options "-fopt-info-omp-all" }
|
||||
{ dg-additional-options "-foffload=-fopt-info-all-omp" } */
|
||||
|
||||
|
|
@ -8,6 +10,15 @@
|
|||
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] }
|
||||
{ dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
|
||||
"WARNING: dg-line var l_dummy defined, but not used". */
|
||||
|
||||
|
||||
#include <assert.h>
|
||||
|
||||
|
||||
|
|
@ -21,9 +32,10 @@ f (void)
|
|||
for (i = 0; i < N; i++)
|
||||
A[i] = -i;
|
||||
|
||||
#pragma acc kernels
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */
|
||||
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
for (i = 0; i < N; i++)
|
||||
A[i] = i;
|
||||
|
||||
|
|
@ -49,15 +61,14 @@ f_data (void)
|
|||
for (i = 0; i < N; i++)
|
||||
A[i] = -i;
|
||||
|
||||
/* See 'declare-vla-kernels-decompose.c'. */
|
||||
#ifdef KERNELS_DECOMPOSE_ICE_HACK
|
||||
(volatile int *) &i;
|
||||
(volatile int *) &N;
|
||||
#endif
|
||||
|
||||
# pragma acc kernels
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */
|
||||
# pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'N' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'N' made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute }
|
||||
{ dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
for (i = 0; i < N; i++)
|
||||
A[i] = i;
|
||||
|
||||
|
|
@ -78,6 +89,3 @@ main ()
|
|||
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
/* { dg-note dummy "" { target n-on-e } } to disable 'prune_notes'. */
|
||||
|
|
|
|||
|
|
@ -24,7 +24,9 @@
|
|||
static int g1;
|
||||
static int g2;
|
||||
|
||||
int main()
|
||||
/* PR100280, etc. */
|
||||
|
||||
static void f1 ()
|
||||
{
|
||||
int a = 0;
|
||||
/*TODO Without making 'a' addressable, for GCN offloading we will not see the expected value copied out. (But it does work for nvptx offloading, strange...) */
|
||||
|
|
@ -153,5 +155,67 @@ int main()
|
|||
assert (g2 == N * (N + 1) / 2);
|
||||
assert (f1 == 2432902008176640000ULL);
|
||||
|
||||
#undef N
|
||||
}
|
||||
|
||||
|
||||
/* PR104086 */
|
||||
|
||||
static void f2 ()
|
||||
{
|
||||
#pragma acc data
|
||||
/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } .-1 } */
|
||||
{
|
||||
int i;
|
||||
|
||||
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i = 1;
|
||||
|
||||
assert (i == 1);
|
||||
|
||||
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
i = -1;
|
||||
|
||||
assert (i == -1);
|
||||
}
|
||||
|
||||
|
||||
int ia[1];
|
||||
|
||||
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'ia' made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
ia[0] = -2;
|
||||
|
||||
assert (ia[0] == -2);
|
||||
|
||||
#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'ia' already made addressable} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute }
|
||||
{ dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
|
||||
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
|
||||
for (int i = 0; i < 100; ++i)
|
||||
++ia[0];
|
||||
|
||||
assert (ia[0] == -2 + 100);
|
||||
}
|
||||
|
||||
|
||||
int main()
|
||||
{
|
||||
f1 ();
|
||||
|
||||
f2 ();
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue