[Bug libgomp/81778] New: libgomp.c/for-5.c failure on nvptx -- illegal memory access

[vries at gcc dot gnu.org]

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

            Bug ID: 81778
           Summary: libgomp.c/for-5.c failure on nvptx -- illegal memory
                    access
           Product: gcc
           Version: unknown
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

At r250889, I noticed libgomp.c/for-5.c fail for nvptx offloading:

Minimized example:
...
extern void abort ();

#pragma omp declare target

#define N 64

int a[N];

#pragma omp end declare target

__attribute__((noinline, noclone)) void
f2_tpf_simd_static32 (void)
{
  unsigned int i;
#pragma omp target parallel for simd schedule(static, 32)
  for (i = N; i > 0 ; i -= 1)
    a[(i - 1)] -= 4;
}

__attribute__((noinline, noclone)) int
test_tpf_simd_static32 (void)
{
  int i;
  for (i = 0; i < N; i++)
    a[i] = i - 25;

#pragma omp target update to(a)

  f2_tpf_simd_static32 ();

#pragma omp target update from(a)

  for (i = 0; i < N; i++)
    if (a[i] != i - 29)
      return 1;

  return 0;
}

int
main ()
{
  if (test_tpf_simd_static32 ())
    abort ();

  return 0;
}
...

...
$ ./install/bin/gcc \
  -fopenmp \
  -B$(pwd -P)/install/offload-nvptx-none/libexec/gcc/x86_64-pc-linux-gnu/8.0.0
\
  -B$(pwd -P)/install/offload-nvptx-none/bin \
  for-5.c \
  -O2
$ ( export LD_LIBRARY_PATH=$(pwd -P)/install/lib64/; ./a.out ; echo $? )

libgomp: cuCtxSynchronize error: an illegal memory access was encountered

libgomp: cuMemFreeHost error: an illegal memory access was encountered

libgomp: device finalization failed
1
...