| Message ID | IA1PR12MB90312B7DB030F838C6AB501BCE912@IA1PR12MB9031.namprd12.prod.outlook.com |
|---|---|
| State | New |
| Headers | show |
| Series | [gimplify.cc] Avoid ICE when passing VLA vector to accelerator | expand |
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)) 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 <prathameshk@nvidia.com> > > Thanks, > Prathamesh > >
> -----Original Message----- > From: Richard Biener <rguenther@suse.de> > Sent: Monday, September 2, 2024 12:47 PM > To: Prathamesh Kulkarni <prathameshk@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 <prathameshk@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 <prathameshk@nvidia.com> > > > > Thanks, > > Prathamesh > > > > > > -- > Richard Biener <rguenther@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 <prathameshk@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);
On Tue, 3 Sep 2024, Prathamesh Kulkarni wrote: > > -----Original Message----- > > From: Richard Biener <rguenther@suse.de> > > Sent: Monday, September 2, 2024 12:47 PM > > To: Prathamesh Kulkarni <prathameshk@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 <prathameshk@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 <prathameshk@nvidia.com> > > > > > > Thanks, > > > Prathamesh > > > > > > > > > > -- > > Richard Biener <rguenther@suse.de> > > SUSE Software Solutions Germany GmbH, > > Frankenstrasse 146, 90461 Nuernberg, Germany; > > GF: Ivo Totev, Andrew McDonald, Werner Knoblich; (HRB 36809, AG > > Nuernberg) >
> -----Original Message----- > From: Richard Biener <rguenther@suse.de> > Sent: Tuesday, September 3, 2024 2:09 PM > To: Prathamesh Kulkarni <prathameshk@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 <rguenther@suse.de> > > > Sent: Monday, September 2, 2024 12:47 PM > > > To: Prathamesh Kulkarni <prathameshk@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 <prathameshk@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 <prathameshk@nvidia.com> > > > > > > > > Thanks, > > > > Prathamesh > > > > > > > > > > > > > > -- > > > Richard Biener <rguenther@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 <rguenther@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 <prathameshk@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);
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 26a216e151d..fb7bd919b54 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7789,6 +7789,11 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) the parameters of the type. */ if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { + /* For now, bail out if a VLA vector is passed to accelerator. */ + if (VECTOR_TYPE_P (TREE_TYPE (decl)) + && POLY_INT_CST_P (DECL_SIZE (decl))) + fatal_error (input_location, + "passing VLA vector to accelerator not implemented"); /* Add the pointer replacement variable as PRIVATE if the variable replacement is private, else FIRSTPRIVATE since we'll need the address of the original variable either for SHARED, or for the
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 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 <prathameshk@nvidia.com> Thanks, Prathamesh [gimplify.cc] Emit an error if VLA vector is passed to accelerator. gcc/ChangeLog: * gimplify.cc (omp_add_variable): Emit an error if VLA vector is passed to accelerator. Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>