On 10-06-19 23:46, Jakub Jelinek wrote:
> Hi!
>
> In PTX, and accepts only .pred, .b16, .b32 and .b64 types, not .u64 etc.
> Fixed thusly, bootstrapped/regtested on x86_64-linux with nvptx-none
> offloading, ok for trunk and 9.2/8.4?
>
LGTM [and to Alexander as mentioned in PR90811 comment 3].
Thanks,
- Tom
> I'll try to figure out tomorrow if we can avoid the overaligned variables in
> this specific testcase, but generally there is always the possibility of
> overaligned vars.
>
> 2019-06-10 Jakub Jelinek <ja...@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.
>
> --- gcc/config/nvptx/nvptx.c.jj 2019-03-11 22:56:55.934666848 +0100
> +++ gcc/config/nvptx/nvptx.c 2019-06-10 15:20:43.154588406 +0200
> @@ -1475,7 +1475,7 @@ nvptx_output_softstack_switch (FILE *fil
> 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)
> --- libgomp/testsuite/libgomp.c/pr90811.c.jj 2019-06-10 15:43:30.464115978
> +0200
> +++ libgomp/testsuite/libgomp.c/pr90811.c 2019-06-10 15:43:37.618003812
> +0200
> @@ -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;
> +}
>
> Jakub
>
