[Bug middle-end/118200] New: note_simd_array_uses crashes in SIMT region with offloading to nvptx

[prathamesh3492 at gcc dot gnu.org via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=118200

            Bug ID: 118200
           Summary: note_simd_array_uses crashes in SIMT region with
                    offloading to nvptx
           Product: gcc
           Version: 15.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: prathamesh3492 at gcc dot gnu.org
  Target Milestone: ---

Hi,
While working on patch to enable SVE with offloading
(https://gcc.gnu.org/pipermail/gcc/2024-December/245265.html),
I came across an issue in note_simd_array_uses, which seems to crash while
processing GOMP internal functions (GOMP_SIMD_VF) in SIMT region. This happens
because in SIMT region, first argument to GOMP internal function is lhs of call
to IFN_GOMP_SIMT_ENTER, rather than default definition, and
note_simd_array_uses crashes at:
ns.simduid = DECL_UID (SSA_NAME_VAR (gimple_call_arg (stmt, 0)));

because the SSA temporary name in this case has SSA_NAME_VAR set to NULL.

This can also be reproduced on trunk with the following test:
#include <stdlib.h>
#define N 1000
int A[N];
int B[N];
int main()
{
  int i;
  int sum = 0;
  #pragma omp target map(sum), map(A), map(B), map (i)
  #pragma omp teams distribute parallel for simd reduction(+:sum)
  for (i = 0; i < N; i++)
    sum += A[i] * B[i];
  return sum;
}

and with following set of options: -O1 -fopenmp -foffload=nvptx-none
-fno-tree-ccp -fno-tree-dce -fno-tree-fre -fno-tree-vrp -fno-tree-forwprop
-fno-tree-copy-prop -fno-tree-dominator-opts

(The rationale behind these set of options is to disable elimination of dead
code that results after omp_device_lower pass).

which results in following ICE:
during GIMPLE pass: vect
dump file: a-foo.c.182t.vect
foo.c: In function ‘main._omp_fn.1’:
foo.c:13:11: internal compiler error: Segmentation fault
   13 |   #pragma omp teams distribute parallel for simd reduction(+:sum)
      |           ^~~
0x244e3e3 internal_error(char const*, ...)
        ../../gcc/gcc/diagnostic-global-context.cc:517
0x103d8bf crash_signal
        ../../gcc/gcc/toplev.cc:322
0x13fe9a4 contains_struct_check(tree_node*, tree_node_structure_enum, char
const*, int, char const*)
        ../../gcc/gcc/tree.h:3810
0x13fe9a4 note_simd_array_uses
        ../../gcc/gcc/tree-vectorizer.cc:421
0x14021db execute
        ../../gcc/gcc/tree-vectorizer.cc:1274

In this case, input to ompdevlow pass is:
  <bb 15> [local count: 118111600]:
  simduid.12_73 = .GOMP_SIMT_ENTER (simduid.12_71(D), 0B, 0B);
  .omp_simt.13_75 = .GOMP_SIMT_ENTER_ALLOC (simduid.12_73);
  _76 = .GOMP_SIMD_VF (simduid.12_73);

which is then folded to:
  <bb 15> [local count: 118111600]:
  simduid.12_73 = simduid.12_71(D);
  .omp_simt.13_75 = 0B;
  _76 = .GOMP_SIMD_VF (simduid.12_73);

and input to vect pass:
  <bb 13> [local count: 118111600]:
  simduid.12_73 = simduid.12_71(D);
  _76 = .GOMP_SIMD_VF (simduid.12_73);

When worked around the issue in note_simd_array_uses, it similarly fails in
adjust_simduid_builtins at:
data.simduid = DECL_UID (SSA_NAME_VAR (arg));

since arg (simduid.12_73) is an argument to one of the GOMP internal functions
and similarly expects it to be default def. 

Thanks,
Prathamesh