public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
@ 2021-05-05 7:33 vries at gcc dot gnu.org
2021-05-05 7:34 ` [Bug target/100428] " vries at gcc dot gnu.org
` (5 more replies)
0 siblings, 6 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-05 7:33 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
Bug ID: 100428
Summary: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-
7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0
-foffload=nvptx-none -O0 execution test
Product: gcc
Version: 12.0
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
Target Milestone: ---
When running libgomp testsuite using GOMP_NVPTX_JIT=-O0, I run into:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
...
In more detail:
...
Execution timeout is: 300
spawn [open ...]^M
reduction-7.exe:
/home/vries/oacc/trunk/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c:312:
v_p_2: Assertion `out[j * 32 + i] == (i + j) * 2' failed.
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
...
Using driver 450.119.03.
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
@ 2021-05-05 7:34 ` vries at gcc dot gnu.org
2021-05-05 11:27 ` vries at gcc dot gnu.org
` (4 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-05 7:34 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Likewise, c++:
...
FAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
...
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
2021-05-05 7:34 ` [Bug target/100428] " vries at gcc dot gnu.org
@ 2021-05-05 11:27 ` vries at gcc dot gnu.org
2021-05-05 11:36 ` vries at gcc dot gnu.org
` (3 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-05 11:27 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Target| |nvptx
--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Minimized to:
...
#include <stdio.h>
#include <assert.h>
int
main (void)
{
unsigned int i, j;
unsigned int out;
#pragma acc parallel num_gangs(1) num_workers(1) vector_length(32) \
copyout(out)
{
for (j = 1; j < 2; j++)
{
unsigned int acc = 0;
#pragma acc loop vector reduction(+:acc)
for (i = 0; i < 32; i++)
acc += 1;
out = acc;
}
}
printf ("out: %u\n", out);
fflush (0);
assert (out == 32);
return 0;
}
...
which gives us:
...
Execution timeout is: 300
spawn [open ...]^M
out: 1
reduction-7.exe: reduction-7.c:27: main: Assertion `out == 32' failed.
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
...
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
2021-05-05 7:34 ` [Bug target/100428] " vries at gcc dot gnu.org
2021-05-05 11:27 ` vries at gcc dot gnu.org
@ 2021-05-05 11:36 ` vries at gcc dot gnu.org
2021-12-10 12:09 ` vries at gcc dot gnu.org
` (2 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-05 11:36 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
The ptx code looks a lot like the cuda reproducer in PR99932 comment 4, so I'm
going to retest this once I get a driver where that one is fixed.
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
` (2 preceding siblings ...)
2021-05-05 11:36 ` vries at gcc dot gnu.org
@ 2021-12-10 12:09 ` vries at gcc dot gnu.org
2022-02-01 18:29 ` cvs-commit at gcc dot gnu.org
2022-02-02 12:21 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-12-10 12:09 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
FTR, reproduces with driver version 470.86 on Quadro M1200 and GeForce GT 1030.
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
` (3 preceding siblings ...)
2021-12-10 12:09 ` vries at gcc dot gnu.org
@ 2022-02-01 18:29 ` cvs-commit at gcc dot gnu.org
2022-02-02 12:21 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: cvs-commit at gcc dot gnu.org @ 2022-02-01 18:29 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
--- Comment #5 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:
https://gcc.gnu.org/g:456de10c549379b74d4858f00d4b8817035a73fc
commit r12-6967-g456de10c549379b74d4858f00d4b8817035a73fc
Author: Tom de Vries <tdevries@suse.de>
Date: Sun Jan 23 06:42:24 2022 +0100
[nvptx] Handle nop in prevent_branch_around_nothing
When running libgomp test-case reduction-7.c on an nvptx accelerator
(T400, driver version 470.86) and GOMP_NVPTX_JIT=-O0, I run into:
...
reduction-7.exe:reduction-7.c:312: v_p_2: \
Assertion `out[j * 32 + i] == (i + j) * 2' failed.
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none \
-O0 execution test
...
During investigation I found ptx code like this:
...
@ %r163 bra $L262;
$L262:
...
There's a known problem with executing this type of code, and a workaround
is
in place to address this: prevent_branch_around_nothing. The workaround
does
not trigger though because it doesn't handle the nop insn.
Fix this by handling the nop insn in prevent_branch_around_nothing.
Tested libgomp on x86_64 with nvptx accelerator.
gcc/ChangeLog:
2022-01-27 Tom de Vries <tdevries@suse.de>
PR target/100428
* config/nvptx/nvptx.cc (prevent_branch_around_nothing): Handle nop
insn.
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Bug target/100428] [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
` (4 preceding siblings ...)
2022-02-01 18:29 ` cvs-commit at gcc dot gnu.org
@ 2022-02-02 12:21 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2022-02-02 12:21 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100428
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Resolution|--- |FIXED
Status|UNCONFIRMED |RESOLVED
Target Milestone|--- |12.0
--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed in gcc-12.
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2022-02-02 12:21 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-05 7:33 [Bug target/100428] New: [nvptx, GOMP_NVPTX_JIT=-O0] FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test vries at gcc dot gnu.org
2021-05-05 7:34 ` [Bug target/100428] " vries at gcc dot gnu.org
2021-05-05 11:27 ` vries at gcc dot gnu.org
2021-05-05 11:36 ` vries at gcc dot gnu.org
2021-12-10 12:09 ` vries at gcc dot gnu.org
2022-02-01 18:29 ` cvs-commit at gcc dot gnu.org
2022-02-02 12:21 ` vries at gcc dot gnu.org
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).