[Bug libgomp/81778] 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

--- Comment #9 from Tom de Vries <vries at gcc dot gnu.org> ---
I ran into this again, and did another round of minimizing.  This time, I added
some buffering around where we write, and check the entire buffer afterwards:
...
$ cat libgomp/testsuite/libgomp.c-c++-common/for-5.c      
/* { dg-additional-options "-std=gnu99" { target c } } */

#include <stdio.h>

#define N 4096
#define MID (N/2)
#define M 4
#pragma omp declare target
int a[N];
#pragma omp end declare target

int
main (void)
{
  int i;

  for (i = 0; i < N; i++)
    a[i] = 0;

#pragma omp target update to(a)

  int s = 1;

#pragma omp target simd
  for (int i = M - 1; i > -1; i -= s)
    a[MID + i] = 1;

#pragma omp target update from(a)

  int error_found = 0;
  for (i = 0; i < N; i++)
    {
      int expected = (MID <= i && i < MID + M) ? 1 : 0;
      if (a[i] == expected)
        continue;

      error_found = 1;
      printf ("Expected %d, got %u at %d\n", expected, a[i], i);
    }

  if (error_found)
    __builtin_abort ();

  return 0;
}
...

Indeed we're writing more than required (for M == 4, we just want the locations
2048..2051):
...
$ LD_LIBRARY_PATH=$(pwd -P)/install/lib64 ./for-5.exe
Expected 0, got 1 at 1955
Expected 0, got 1 at 1986
Expected 0, got 1 at 1987
Expected 0, got 1 at 2017
Expected 0, got 1 at 2018
Expected 0, got 1 at 2019
Aborted (core dumped)
...

Fails for M >= 2.