On Mon, 19 Aug 2024, Prathamesh Kulkarni wrote:
> Hi Richard,
> As mentioned in RFC email, for the following test:
>
> int main()
> {
> long c[4];
> #pragma omp target map(c)
> c[0] = 0;
> return 0;
> }
>
> Compiling for AArch64 host with -O2 -fopenmp -foffload=nvptx-none results in:
> lto1: fatal error: nvptx-none - 256-bit integer numbers unsupported (mode
> 'OI') compilation terminated.
> nvptx mkoffload: fatal error:
> ../install/bin/aarch64-unknown-linux-gnu-accel-nvptx-none-gcc returned 1 exit
> status compilation terminated.
>
> This happens because AArch64 uses OImode for ARRAY_TYPE whose size fits
> 256-bits, which is not supported on nvptx, and thus
> emits the above diagnostic.
>
> Following your suggestion, the attached patch streams out VOIDmode from host
> for TYPE_MODE and DECL_MODE for aggregate types
> with offloading enabled, and while streaming-in on accel side, it recomputes
> TYPE_MODE and DECL_MODE, which fixes the issue.
> Patch survives AArch64->nvptx offload testing for libgomp and bootstrap+test
> on aarch64-linux-gnu.
>
> Does the patch look in the right direction ?
+/* Compute TYPE_MODE for TYPE (which is ARRAY_TYPE). */
+
+void compute_array_mode (tree type)
+{
newline after 'void'
Otherwise LGTM, please leave time for others to comment though.
Thanks,
Richard.
> Signed-off-by: Prathamesh Kulkarni <prathame...@nvidia.com>
>
> Thanks,
> Prathamesh
>
--
Richard Biener <rguent...@suse.de>
SUSE Software Solutions Germany GmbH,
Frankenstrasse 146, 90461 Nuernberg, Germany;
GF: Ivo Totev, Andrew McDonald, Werner Knoblich; (HRB 36809, AG Nuernberg)
