mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-04 06:30:25 +08:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source 
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
This commit is contained in:
Jakub Jelinek 2019-06-11 18:40:10 +02:00 committed by Jakub Jelinek
parent d37c29f942
commit ce9c4ec3c5
4 changed files with 41 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
gcc/ChangeLog
View File

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

2
gcc/config/nvptx/nvptx.c
View File

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

5
libgomp/ChangeLog
View File

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

29
libgomp/testsuite/libgomp.c/pr90811.c Normal file
View File

@ -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;
}