Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086]

Message ID 87ilsjmjld.fsf@euler.schwinge.homeip.net
State New
Headers
Series Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086] |

Commit Message

Thomas Schwinge March 12, 2022, 12:38 p.m. UTC
  Hi!

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'.
> +   { 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]",
see attached.


Grüße
 Thomas


> --- /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;


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

Patch

From 9781ae3a254a8c17ef4ffa70f21ed1728ff3c707 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 18 Jan 2022 17:22:14 +0100
Subject: [PATCH] Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c'
 [PR104086]

..., currently XFAILed with 'dg-ice', as it runs into
'gcc/omp-low.cc:lower_omp_target':

    13125			else if (is_gimple_reg (var))
    13126			  {
    13127			    gcc_assert (offloaded);

This means, the recent PR100280 etc. changes are still not sufficient.

	gcc/testsuite/
	PR middle-end/104086
	* c-c++-common/goacc/kernels-decompose-pr104086-1.c: New file.
---
 .../goacc/kernels-decompose-pr104086-1.c      | 25 +++++++++++++++++++
 1 file changed, 25 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c
new file mode 100644
index 00000000000..eab10cf6c72
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c
@@ -0,0 +1,25 @@ 
+/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c'.  */
+
+/* { dg-additional-options "-fchecking" }
+   { dg-ice TODO }
+   { dg-prune-output {during GIMPLE pass: omplower} } */
+
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+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 } */
+  {
+    int i;
+
+#pragma acc kernels
+    /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+    i = 0;
+  }
+}
-- 
2.34.1