diff options
| author | Jakub Jelinek <jakub@redhat.com> | 2019-06-11 18:40:10 +0200 |
|---|---|---|
| committer | Jakub Jelinek <jakub@gcc.gnu.org> | 2019-06-11 18:40:10 +0200 |
| commit | ce9c4ec3c5e70023b27e49c2fb45eafaddec9a14 (patch) | |
| tree | 263460c3e2e396e21941f20240ba4c778b86c656 | |
| parent | d37c29f94234b8b62334f895b69c992bea6a6f10 (diff) | |
| download | gcc-ce9c4ec3c5e70023b27e49c2fb45eafaddec9a14.zip gcc-ce9c4ec3c5e70023b27e49c2fb45eafaddec9a14.tar.gz gcc-ce9c4ec3c5e70023b27e49c2fb45eafaddec9a14.tar.bz2 | |
re PR target/90811 ([nvptx] ptxas error on OpenMP offloaded code)
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: r272161
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 2 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 5 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/pr90811.c | 29 |
4 files changed, 41 insertions, 1 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index e4776e0..95cdab7 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +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-06-11 Marc Glisse <marc.glisse@inria.fr> * match.pd (X/[ex]4<Y/[ex]4): Handle conversions. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index a28099a..1986e79 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -1475,7 +1475,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 2ea937e..3487880 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2019-06-11 Jakub Jelinek <jakub@redhat.com> + + PR target/90811 + * testsuite/libgomp.c/pr90811.c: New test. + 2019-06-05 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c++/lastprivate-conditional-1.C: New test. diff --git a/libgomp/testsuite/libgomp.c/pr90811.c b/libgomp/testsuite/libgomp.c/pr90811.c new file mode 100644 index 0000000..25b7d78 --- /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; +} |