[Bug middle-end/117003] New: pr104783.c is miscompiled with offloading and results in segmentation fault during host-only execution for -O1 and above

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

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

            Bug ID: 117003
           Summary: pr104783.c is miscompiled with offloading and results
                    in segmentation fault during host-only execution for
                    -O1 and above
           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,
The following test (libgomp/pr104783.c):

int
main (void)
{
  unsigned val = 0;

#pragma omp target map(tofrom: val)
#pragma omp simd
  for (int i = 0 ; i < 1 ; i++)
    {
#pragma omp atomic update
      val = val + 1;
    }

  if (val != 1)
    __builtin_abort ();

  return 0;
}

Compiling with -O1 -fopenmp -foffload=nvptx-none and forcing host-only
execution results in segfault.

The issue here seems to be during omp lowering.

                  D.4569 = .GOMP_USE_SIMT ();
                  if (D.4569 != 0) goto <D.4570>; else goto <D.4571>;
                  <D.4570>:
                  {
                    void * simduid.2;
                    void * .omp_simt.3;
                    int i;

                    simduid.2 = .GOMP_SIMT_ENTER (simduid.2);
                    .omp_simt.3 = .GOMP_SIMT_ENTER_ALLOC (simduid.2);
                    #pragma omp simd _simduid_(simduid.2) _simt_ linear(i:1)
                    for (i = 0; i < 1; i = i + 1)
                    D.4577 = .omp_data_i->val;
                    #pragma omp atomic_load relaxed
                      D.4558 = *D.4577
                    D.4559 = D.4558 + 1;
                    #pragma omp atomic_store relaxed (D.4559)
                    #pragma omp continue (i, i)
                    .GOMP_SIMT_EXIT (.omp_simt.3);
                    #pragma omp return(nowait)
                  }
                 goto <D.4572>;
                 <D.4571>:
                  {
                    int i;

                    #pragma omp simd linear(i:1)
                    for (i = 0; i < 1; i = i + 1)
                    #pragma omp atomic_load relaxed
                      D.4558 = *&*D.4577
                    D.4559 = D.4558 + 1;
                    #pragma omp atomic_store relaxed (D.4559)
                    #pragma omp continue (i, i)
                    #pragma omp return(nowait)
                  }
                 <D.4572>:
                 ...

In the following stmt in simd code-path:
  #pragma omp atomic_load relaxed
    D.4558 = *&*D.4577

D.4577 is uninitialized. D.4577 is initialized in the sibling
if-block containing simt code-path:
D.4577 = .omp_data_i->val;

which doesn't reach the use in atomic_load relaxed stmt,
and gets expanded to following in ompexp dump:
  <bb 10> :
  __atomic_fetch_add_4 (D.4590, 1, 0);

and thus we end up passing uninitialized pointer to __atomic_fetch_add_4, which
results in segfault.

While in the corresponding simt code-path, it's initialized correctly:
  <bb 5> :
  D.4590 = .omp_data_i->val;
  __atomic_fetch_add_4 (D.4590, 1, 0);

Thanks,
Prathamesh