> -----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 ?
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)
Avoid ICE when passing VLA vector to accelerator.
gcc/ChangeLog:
* gimplify.cc (omp_add_variable): Check if decl size is not
POLY_INT_CST.
(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 081d69bce05..fd3a451f4bc 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7799,7 +7799,9 @@ 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)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+ && !POLY_INT_CST_P (DECL_SIZE (decl)))
{
/* Add the pointer replacement variable as PRIVATE if the variable
replacement is private, else FIRSTPRIVATE since we'll need the
@@ -14414,6 +14416,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_CST_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..466541d4255 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1664,7 +1664,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
- && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+ && !POLY_INT_CST_P (DECL_SIZE (decl)))
{
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (INDIRECT_REF_P (decl2));
@@ -1906,7 +1907,8 @@ 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)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+ && !POLY_INT_CST_P (DECL_SIZE (decl)))
{
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (INDIRECT_REF_P (decl2));
@@ -12750,7 +12752,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
if (DECL_SIZE (var)
- && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST
+ && !POLY_INT_CST_P (DECL_SIZE (var)))
{
tree var2 = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
@@ -13077,7 +13080,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
else
{
if (DECL_SIZE (ovar)
- && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+ && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST
+ && !POLY_INT_CST_P (DECL_SIZE (ovar)))
{
tree ovar2 = DECL_VALUE_EXPR (ovar);
gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
