[gimplify.cc] Avoid ICE when passing VLA vector to accelerator

[Prathamesh Kulkarni]

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 <prathame...@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 <prathame...@nvidia.com>

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