On Tue, 3 Sep 2024, Prathamesh Kulkarni wrote:
> > -----Original Message-----
> > From: Richard Biener <rguent...@suse.de>
> > Sent: Monday, September 2, 2024 12:47 PM
> > To: Prathamesh Kulkarni <prathame...@nvidia.com>
> > Cc: gcc-patches@gcc.gnu.org
> > Subject: Re: [gimplify.cc] Avoid ICE when passing VLA vector to
> > accelerator
> >
> > External email: Use caution opening links or attachments
> >
> >
> > On Sun, 1 Sep 2024, Prathamesh Kulkarni wrote:
> >
> > > Hi,
> > > For the following test:
> > > #include <arm_sve.h>
> > >
> > > int main()
> > > {
> > > svint32_t x;
> > > #pragma omp target map(x)
> > > x;
> > > return 0;
> > > }
> > >
> > > compiling with -fopenmp -foffload=nvptx-none results in following
> > ICE:
> > >
> > > t_sve.c: In function 'main':
> > > t_sve.c:6:11: internal compiler error: Segmentation fault
> > > 6 | #pragma omp target map(x)
> > > | ^~~
> > > 0x228ed13 internal_error(char const*, ...)
> > > ../../gcc/gcc/diagnostic-global-context.cc:491
> > > 0xfcf68f crash_signal
> > > ../../gcc/gcc/toplev.cc:321
> > > 0xc17d9c omp_add_variable
> > > ../../gcc/gcc/gimplify.cc:7811
> >
> > that's not on trunk head? Anyway, I think that instead
> >
> > /* When adding a variable-sized variable, we have to handle all
> > sorts
> > of additional bits of data: the pointer replacement variable, and
> > the parameters of the type. */
> > if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
> >
> > should instead be checking for !POLY_INT_CST_P (DECl_SIZE (decl))
> Hi Richard,
> Thanks for the suggestions. The attached patch adds !POLY_INT_CST_P check in
> omp_add_variable
> (and few more places where it segfaulted), but keeps TREE_CODE (DECL_SIZE
> (decl)) != INTEGER_CST check to
> avoid above ICE with -msve-vector-bits= option.
>
> The test now fails with:
> lto1: fatal error: degree of 'poly_int' exceeds 'NUM_POLY_INT_COEFFS' (1)
> compilation terminated.
> nvptx mkoffload: fatal error:
> ../install/bin/aarch64-unknown-linux-gnu-accel-nvptx-none-gcc returned 1 exit
> status
> compilation terminated.
>
> Which looks reasonable IMO, since we don't yet fully support streaming of
> poly_ints
> (and compiles OK when length is set with -msve-vector-bits= option).
>
> Bootstrap+test in progress on aarch64-linux-gnu.
> Does the patch look OK ?
Please use use !poly_int_tree_p which checks for both
INTEGER_CST and POLY_INT_CST_P.
OK with that change.
Richard.
>
> Signed-off-by: Prathamesh Kulkarni <prathame...@nvidia.com>
>
> Thanks,
> Prathamesh
> >
> > Richard.
> >
> >
> > > 0xc17d9c omp_add_variable
> > > ../../gcc/gcc/gimplify.cc:7752 0xc4176b
> > > gimplify_scan_omp_clauses
> > > ../../gcc/gcc/gimplify.cc:12881
> > > 0xc46d53 gimplify_omp_workshare
> > > ../../gcc/gcc/gimplify.cc:17139
> > > 0xc23383 gimplify_expr(tree_node**, gimple**, gimple**, bool
> > (*)(tree_node*), int)
> > > ../../gcc/gcc/gimplify.cc:18668
> > > 0xc27f53 gimplify_stmt(tree_node**, gimple**)
> > > ../../gcc/gcc/gimplify.cc:7646
> > > 0xc24ef7 gimplify_statement_list
> > > ../../gcc/gcc/gimplify.cc:2250
> > > 0xc24ef7 gimplify_expr(tree_node**, gimple**, gimple**, bool
> > (*)(tree_node*), int)
> > > ../../gcc/gcc/gimplify.cc:18565
> > > 0xc27f53 gimplify_stmt(tree_node**, gimple**)
> > > ../../gcc/gcc/gimplify.cc:7646
> > > 0xc289d3 gimplify_bind_expr
> > > ../../gcc/gcc/gimplify.cc:1642 0xc24b9b
> > > gimplify_expr(tree_node**, gimple**, gimple**, bool (*)(tree_node*),
> > int)
> > > ../../gcc/gcc/gimplify.cc:18315
> > > 0xc27f53 gimplify_stmt(tree_node**, gimple**)
> > > ../../gcc/gcc/gimplify.cc:7646
> > > 0xc24ef7 gimplify_statement_list
> > > ../../gcc/gcc/gimplify.cc:2250
> > > 0xc24ef7 gimplify_expr(tree_node**, gimple**, gimple**, bool
> > (*)(tree_node*), int)
> > > ../../gcc/gcc/gimplify.cc:18565
> > > 0xc27f53 gimplify_stmt(tree_node**, gimple**)
> > > ../../gcc/gcc/gimplify.cc:7646 0xc2aadb
> > > gimplify_body(tree_node*, bool)
> > > ../../gcc/gcc/gimplify.cc:19393 0xc2b05f
> > > gimplify_function_tree(tree_node*)
> > > ../../gcc/gcc/gimplify.cc:19594 0xa0e47f
> > > cgraph_node::analyze()
> > > ../../gcc/gcc/cgraphunit.cc:687
> > >
> > > The attached patch fixes the issue by checking if variable is VLA
> > > vector, and emits an error in that case since no accel currently
> > supports VLA vectors.
> > > Does the patch look OK ?
> > >
> > > 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)
>
--
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)
