From: Jakub Jelinek <jakub@redhat.com>
Date: Fri, 30 Aug 2019 12:43:04 +0000 (+0200)
Subject: backport: re PR target/90811 ([nvptx] ptxas error on OpenMP offloaded code)
X-Git-Tag: releases/gcc-7.5.0~208
X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=57bc4de20a60befb462b0b91850eeba186515bfa;p=thirdparty%2Fgcc.git

backport: re PR target/90811 ([nvptx] ptxas error on OpenMP offloaded code)

	Backported from mainline
	2019-06-11  Jakub Jelinek  <jakub@redhat.com>

	PR target/90811
	* config/nvptx/nvptx.c (nvptx_output_softstack_switch): Use and.b%d
	instead of and.u%d.

	* testsuite/libgomp.c/pr90811.c: New test.

From-SVN: r275155
---

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 6a9190f8b422..2e41dffc42c4 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,6 +1,12 @@
 2019-08-30  Jakub Jelinek  <jakub@redhat.com>
 
 	Backported from mainline
+	2019-06-11  Jakub Jelinek  <jakub@redhat.com>
+
+	PR target/90811
+	* config/nvptx/nvptx.c (nvptx_output_softstack_switch): Use and.b%d
+	instead of and.u%d.
+
 	2019-05-29  Jakub Jelinek  <jakub@redhat.com>
 
 	PR fortran/90329
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index af2c4a23dfee..74bce3ec8d0f 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -1354,7 +1354,7 @@ nvptx_output_softstack_switch (FILE *file, bool entering,
       fputs (";\n", file);
       if (!CONST_INT_P (size) || UINTVAL (align) > GET_MODE_SIZE (DImode))
 	fprintf (file,
-		 "\t\tand.u%d %%r%d, %%r%d, -" HOST_WIDE_INT_PRINT_DEC ";\n",
+		 "\t\tand.b%d %%r%d, %%r%d, -" HOST_WIDE_INT_PRINT_DEC ";\n",
 		 bits, regno, regno, UINTVAL (align));
     }
   if (cfun->machine->has_softstack)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 946ac0e99f73..4b2d8ff8d490 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,6 +1,11 @@
 2019-08-30  Jakub Jelinek  <jakub@redhat.com>
 
 	Backported from mainline
+	2019-06-11  Jakub Jelinek  <jakub@redhat.com>
+
+	PR target/90811
+	* testsuite/libgomp.c/pr90811.c: New test.
+
 	2019-01-28  Jakub Jelinek  <jakub@redhat.com>
 
 	PR middle-end/89002
diff --git a/libgomp/testsuite/libgomp.c/pr90811.c b/libgomp/testsuite/libgomp.c/pr90811.c
new file mode 100644
index 000000000000..25b7d78c0eac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr90811.c
@@ -0,0 +1,29 @@
+/* PR target/90811 */
+
+int
+main ()
+{
+  long long a[100], b[100];
+  int i;
+  for (i = 0; i < 100; i++)
+    {
+      a[i] = i;
+      b[i] = i % 10;
+    }
+  #pragma omp target teams distribute parallel for simd map(tofrom: a[:100], b[:100])
+  for (i = 0; i < 100; i++)
+    {
+      long long c = 0;
+      const long long d[] = { 1, 3, 5, 7, 9 };
+      for (int j = 4; j >= 0; j--)
+         c = d[j] + b[i] * c;
+      a[i] += c;
+    }
+  for (i = 0; i < 100; i++)
+    {
+      const long long r[] = { 1, 26, 229, 976, 2849, 6646, 13381, 24284, 40801, 64594 };
+      if (a[i] != r[i % 10] + (i / 10 * 10))
+	__builtin_abort ();
+    }
+  return 0;
+}