> -----Original Message-----
> From: Richard Biener <rguent...@suse.de>
> Sent: Tuesday, September 3, 2024 2:09 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 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.
Thanks, I have committed the attached patch in:
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=ae88e91938af364ef5613e5461b12b484b578bc5
after verifying it passes bootstrap+test on aarch64-linux-gnu and survives
libgomp tests
for Aarch64/nvptx offloading.
Thanks,
Prathamesh
>
> 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)
Avoid ICE when passing VLA vector to accelerator.
gcc/ChangeLog:
* gimplify.cc (omp_add_variable): Check if decl size is not
poly_int_tree_p.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.cc (scan_sharing_clauses): Likewise.
(lower_omp_target): Likewise.
Signed-off-by: Prathamesh Kulkarni <prathame...@nvidia.com>
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 9300138aa0c..ceb53e5d5bb 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7799,7 +7799,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree
decl, unsigned int flags)
/* 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)
+ if (DECL_SIZE (decl) && !poly_int_tree_p (DECL_SIZE (decl)))
{
/* Add the pointer replacement variable as PRIVATE if the variable
replacement is private, else FIRSTPRIVATE since we'll need the
@@ -14413,7 +14413,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
}
}
else if (DECL_SIZE (decl)
- && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+ && !poly_int_tree_p (DECL_SIZE (decl))
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 4d003f42098..241f79e34a9 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1664,7 +1664,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
- && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ && !poly_int_tree_p (DECL_SIZE (decl)))
{
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (INDIRECT_REF_P (decl2));
@@ -1906,7 +1906,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
= remap_type (TREE_TYPE (decl), &ctx->cb);
}
else if (DECL_SIZE (decl)
- && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ && !poly_int_tree_p (DECL_SIZE (decl)))
{
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (INDIRECT_REF_P (decl2));
@@ -12750,7 +12750,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
if (DECL_SIZE (var)
- && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ && !poly_int_tree_p (DECL_SIZE (var)))
{
tree var2 = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
@@ -13077,7 +13077,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
else
{
if (DECL_SIZE (ovar)
- && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+ && !poly_int_tree_p (DECL_SIZE (ovar)))
{
tree ovar2 = DECL_VALUE_EXPR (ovar);
gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
