* [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access
[not found] <bug-81778-4@http.gcc.gnu.org/bugzilla/>
@ 2020-09-25 8:54 ` burnus at gcc dot gnu.org
2020-10-01 13:36 ` vries at gcc dot gnu.org
` (3 subsequent siblings)
4 siblings, 0 replies; 5+ messages in thread
From: burnus at gcc dot gnu.org @ 2020-09-25 8:54 UTC (permalink / raw)
To: gcc-bugs
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
^ permalink raw reply [flat|nested] 5+ messages in thread* [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access
[not found] <bug-81778-4@http.gcc.gnu.org/bugzilla/>
2020-09-25 8:54 ` [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access burnus at gcc dot gnu.org
@ 2020-10-01 13:36 ` vries at gcc dot gnu.org
2020-10-01 23:28 ` vries at gcc dot gnu.org
` (2 subsequent siblings)
4 siblings, 0 replies; 5+ messages in thread
From: vries at gcc dot gnu.org @ 2020-10-01 13:36 UTC (permalink / raw)
To: gcc-bugs
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.
^ permalink raw reply [flat|nested] 5+ messages in thread* [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access
[not found] <bug-81778-4@http.gcc.gnu.org/bugzilla/>
2020-09-25 8:54 ` [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access burnus at gcc dot gnu.org
2020-10-01 13:36 ` vries at gcc dot gnu.org
@ 2020-10-01 23:28 ` vries at gcc dot gnu.org
2020-10-02 15:05 ` burnus at gcc dot gnu.org
2021-04-29 12:40 ` vries at gcc dot gnu.org
4 siblings, 0 replies; 5+ messages in thread
From: vries at gcc dot gnu.org @ 2020-10-01 23:28 UTC (permalink / raw)
To: gcc-bugs
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. */
...
^ permalink raw reply [flat|nested] 5+ messages in thread* [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access
[not found] <bug-81778-4@http.gcc.gnu.org/bugzilla/>
` (2 preceding siblings ...)
2020-10-01 23:28 ` vries at gcc dot gnu.org
@ 2020-10-02 15:05 ` burnus at gcc dot gnu.org
2021-04-29 12:40 ` vries at gcc dot gnu.org
4 siblings, 0 replies; 5+ messages in thread
From: burnus at gcc dot gnu.org @ 2020-10-02 15:05 UTC (permalink / raw)
To: gcc-bugs
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
^ permalink raw reply [flat|nested] 5+ messages in thread* [Bug libgomp/81778] libgomp.c/for-5.c failure on nvptx -- illegal memory access
[not found] <bug-81778-4@http.gcc.gnu.org/bugzilla/>
` (3 preceding siblings ...)
2020-10-02 15:05 ` burnus at gcc dot gnu.org
@ 2021-04-29 12:40 ` vries at gcc dot gnu.org
4 siblings, 0 replies; 5+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-29 12:40 UTC (permalink / raw)
To: gcc-bugs
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
^ permalink raw reply [flat|nested] 5+ messages in thread