[committed,nvptx] Handle assignment to gang-level reduction variable

Message ID 20190115101350.GA15881@delia
State New
Headers show
Series
  • [committed,nvptx] Handle assignment to gang-level reduction variable
Related show

Commit Message

Tom de Vries Jan. 15, 2019, 10:13 a.m.
Hi,

this fixes an ICE when handling an assignment to a gang-level reduction
variable.

Committed to trunk.

Thanks,
- Tom

[nvptx] Handle assignment to gang-level reduction variable

2019-01-15  Tom de Vries  <tdevries@suse.de>

	PR target/80547
	* config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Handle
	lhs == NULL_TREE for gang-level reduction.

	* testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c:
	New test.

---
 gcc/config/nvptx/nvptx.c                                 |  3 ++-
 .../gang-reduction-var-assignment.c                      | 16 ++++++++++++++++
 2 files changed, 18 insertions(+), 1 deletion(-)

Comments

Thomas Schwinge June 8, 2021, 9:14 a.m. | #1
Hi!

On 2019-01-15T11:13:51+0100, Tom de Vries <tdevries@suse.de> wrote:
> this fixes an ICE when handling an assignment to a gang-level reduction

> variable.

>

> Committed to trunk.


>       PR target/80547


Pushed "Revert PR80547 workaround in
'libgomp.oacc-c-c++-common/parallel-dims.c'" to master branch in commit
0886426f5f543e813c1a61e18da6616caf377dfc, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
From 0886426f5f543e813c1a61e18da6616caf377dfc Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>

Date: Mon, 7 Jun 2021 16:17:06 +0200
Subject: [PATCH] Revert PR80547 workaround in
 'libgomp.oacc-c-c++-common/parallel-dims.c'

This problem has been fixed long ago, in r267934 (commit
d41d952c9bbdffe6fd2badc9c4f2c18d241ce412) "[nvptx] Handle assignment to
gang-level reduction variable".

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Revert
	PR80547 workaround.
---
 .../libgomp.oacc-c-c++-common/parallel-dims.c       | 13 -------------
 1 file changed, 13 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index ef3dfda5fa5..c7412c2ef3a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -72,22 +72,9 @@ int main ()
       gangs_actual = 1;
       for (int i = 100 * gangs_actual; i > -100 * gangs_actual; --i)
 	{
-	  /* <https://gcc.gnu.org/PR80547>.  */
-#if 0
 	  gangs_min = gangs_max = acc_gang ();
 	  workers_min = workers_max = acc_worker ();
 	  vectors_min = vectors_max = acc_vector ();
-#else
-	  int gangs = acc_gang ();
-	  gangs_min = (gangs_min < gangs) ? gangs_min : gangs;
-	  gangs_max = (gangs_max > gangs) ? gangs_max : gangs;
-	  int workers = acc_worker ();
-	  workers_min = (workers_min < workers) ? workers_min : workers;
-	  workers_max = (workers_max > workers) ? workers_max : workers;
-	  int vectors = acc_vector ();
-	  vectors_min = (vectors_min < vectors) ? vectors_min : vectors;
-	  vectors_max = (vectors_max > vectors) ? vectors_max : vectors;
-#endif
 	}
     }
     if (gangs_actual != 1)
-- 
2.30.2

Patch

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 03c0f82f4a2..23459e1c6f4 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -6242,7 +6242,8 @@  nvptx_goacc_reduction_init (gcall *call, offload_attrs *oa)
 	    init = var;
 	}
 
-      gimplify_assign (lhs, init, &seq);
+      if (lhs != NULL_TREE)
+	gimplify_assign (lhs, init, &seq);
     }
 
   pop_gimplify_context (NULL);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c
new file mode 100644
index 00000000000..05f58a4bddf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c
@@ -0,0 +1,16 @@ 
+/* { dg-xfail-run-if "PR88852" { openacc_host_selected } } */
+
+int
+main (void)
+{
+  int x = 123;
+
+#pragma acc parallel num_gangs(1) reduction (+: x)
+  {
+    x = 23;
+  }
+  if (x != 146)
+    __builtin_abort();
+
+  return 0;
+}