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

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

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |burnus at gcc dot gnu.org

--- Comment #8 from Tobias Burnus <burnus at gcc dot gnu.org> ---
See also PR95654 + PR97203, and PR80053

[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.

[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 #10 from Tom de Vries <vries at gcc dot gnu.org> ---
Tentative patch:
...
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 99cb4f9dda4..034de497390 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -6333,6 +6333,8 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
       /* Collapsed loops not handled for SIMT yet: limit to one lane only.  */
       if (fd->collapse > 1)
        simt_maxlane = build_one_cst (unsigned_type_node);
+      else if (TREE_CODE (fd->loops[0].step) != INTEGER_CST)
+       simt_maxlane = build_one_cst (unsigned_type_node);
       else if (safelen_int < omp_max_simt_vf ())
        simt_maxlane = build_int_cst (unsigned_type_node, safelen_int);
       tree vf
@@ -6636,6 +6638,10 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
       else
        t = fold_build2 (PLUS_EXPR, type, fd->loop.v, step);
       expand_omp_build_assign (&gsi, fd->loop.v, t);
+      /* The alternative IV needs to to be updated as well, but isn't
+        currently.  Assert that we fall back to simt_maxlane == 1.  */
+      gcc_assert (altv == NULL_TREE
+                 || tree_int_cst_equal (simt_maxlane, integer_one_node));
     }

   /* Remove GIMPLE_OMP_RETURN.  */
...

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

[burnus at gcc dot gnu.org]

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

--- Comment #11 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Cross ref: the submitted patch is at
https://gcc.gnu.org/pipermail/gcc-patches/2020-October/555352.html

[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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
   Target Milestone|---                         |12.0
             Status|NEW                         |RESOLVED

--- Comment #12 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed in
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=fc14ff611181c274584c7963bc597a6ca50c20a1