OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086]
Commit Message
Hi!
On 2022-03-12T13:38:38+0100, I wrote:
> On 2020-11-13T23:22:30+0100, I wrote:
>> On 2019-02-01T00:59:30+0100, I wrote:
>>> I've just pushed the attached nine patches to openacc-gcc-8-branch:
>>> OpenACC 'kernels' construct changes: splitting of the construct into
>>> several regions.
>>
>> Now, slightly more polished, I've pushed to master branch a variant of
>> most of these patches combined in commit
>> e898ce7997733c29dcab9c3c62ca102c7f9fa6eb "Decompose OpenACC 'kernels'
>> constructs into parts, a sequence of compute constructs", see attached.
>>
>>> There's more work to be done there, and we're aware of a number of TODO
>>> items, but nevertheless: it's a good first step.
>>
>> That's still the case... :-)
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c
>> @@ -0,0 +1,8 @@
>> +/* { dg-additional-options "-fopenacc-kernels=decompose" } */
>> +/* Hopefully, this is the same issue as '../../../gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c'.
(Related, but not the same.)
>> + { dg-ice "TODO" }
>> + TODO { dg-prune-output "during GIMPLE pass: omplower" }
>> + TODO { dg-do link } */
>> +
>> +#undef KERNELS_DECOMPOSE_ICE_HACK
>> +#include "declare-vla.c"
>
> Arseny had later reduced that, and filed <https://gcc.gnu.org/PR104086>.
> To document the status quo, pushed to master branch
> commit 9781ae3a254a8c17ef4ffa70f21ed1728ff3c707
> "Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086]"
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c
>> @@ -0,0 +1,6 @@
>> +/* { dg-additional-options "-fopenacc-kernels=decompose" } */
>> +
>> +/* See also 'declare-vla-kernels-decompose-ice-1.c'. */
>> +
>> +#define KERNELS_DECOMPOSE_ICE_HACK
>> +#include "declare-vla.c"
>> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
>> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
>> @@ -38,6 +38,12 @@ 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
>> for (i = 0; i < N; i++)
>> A[i] = i;
Pushed to master branch commit 337ed336d7dd83526891bdb436f0bfe9e351f69d
"OpenACC 'kernels' decomposition: Mark variables used in 'present'
clauses as addressable [PR100280, PR104086]", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
From 337ed336d7dd83526891bdb436f0bfe9e351f69d Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu, 17 Feb 2022 14:18:57 +0100
Subject: [PATCH] OpenACC 'kernels' decomposition: Mark variables used in
'present' clauses as addressable [PR100280, PR104086]
... like in recent commit 9b32c1669aad5459dd053424f9967011348add83
"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.
---
gcc/omp-low.cc | 27 +++++---
gcc/omp-oacc-kernels-decompose.cc | 32 +++++++++
.../goacc/kernels-decompose-pr104086-1.c | 37 +++++++++--
.../declare-vla-kernels-decompose-ice-1.c | 22 -------
.../declare-vla-kernels-decompose.c | 29 --------
.../libgomp.oacc-c-c++-common/declare-vla.c | 38 ++++++-----
.../kernels-decompose-1.c | 66 ++++++++++++++++++-
7 files changed, 168 insertions(+), 83 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c
@@ -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;
+}
deleted file mode 100644
@@ -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 } */
deleted file mode 100644
@@ -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;
}
--
2.34.1