* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
@ 2021-05-10 10:44 ` vries at gcc dot gnu.org
2021-05-10 13:04 ` burnus at gcc dot gnu.org
` (14 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-10 10:44 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Can you post a minimal version?
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
2021-05-10 10:44 ` [Bug target/100497] " vries at gcc dot gnu.org
@ 2021-05-10 13:04 ` burnus at gcc dot gnu.org
2021-05-10 15:04 ` vries at gcc dot gnu.org
` (13 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-05-10 13:04 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #1)
> Can you post a minimal version?
See below, fails with 4 systems, works with 3 others.
#define N 1024
int
main ()
{
_Complex float orfc = 0;
// _Complex int orfc = 0; // variant; uses a GNU extension
int rf[N] = {}; // Can also be _Complex (float/int) as in the orig example
rf[5] = 1;
#pragma omp target teams distribute parallel for reduction(||: orfc)
map(orfc)
for (int i=0; i < N; ++i)
orfc = orfc || rf[i];
if (__real__ orfc != 1.0) // if it fails, it is 0.0 + 0.0i
__builtin_abort ();
return 0;
}
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
2021-05-10 10:44 ` [Bug target/100497] " vries at gcc dot gnu.org
2021-05-10 13:04 ` burnus at gcc dot gnu.org
@ 2021-05-10 15:04 ` vries at gcc dot gnu.org
2021-05-10 15:07 ` vries at gcc dot gnu.org
` (12 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-10 15:04 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
Doesn't fail for me unfortunately.
I've tried with GOMP_NVPTX_JIT=-O0..-O4, no luck.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (2 preceding siblings ...)
2021-05-10 15:04 ` vries at gcc dot gnu.org
@ 2021-05-10 15:07 ` vries at gcc dot gnu.org
2021-05-10 15:08 ` vries at gcc dot gnu.org
` (11 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-10 15:07 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #2)
> See below, fails with 4 systems, works with 3 others.
Can anything be deduced from driver versions?
Or card architecture?
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (3 preceding siblings ...)
2021-05-10 15:07 ` vries at gcc dot gnu.org
@ 2021-05-10 15:08 ` vries at gcc dot gnu.org
2021-05-10 16:21 ` burnus at gcc dot gnu.org
` (10 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-10 15:08 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
Does it pass with GOMP_NVPTX_JIT=-O0 ?
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (4 preceding siblings ...)
2021-05-10 15:08 ` vries at gcc dot gnu.org
@ 2021-05-10 16:21 ` burnus at gcc dot gnu.org
2021-05-10 17:37 ` vries at gcc dot gnu.org
` (9 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-05-10 16:21 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #6 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #5)
> Does it pass with GOMP_NVPTX_JIT=-O0 ?
This does not seem to help either way.
(In reply to Tom de Vries from comment #4)
> Can anything be deduced from driver versions?
> Or card architecture?
Maybe. The driver version does not seem to matter – but the card architecture
could:
* Working
Tesla K40c | NVIDIA-SMI 346.46 Driver Version: 346.46
Tesla K20c | NVIDIA-SMI 455.38 Driver Version: 455.38 CUDA Version: 11.1
Tesla K80 | NVIDIA-SMI 440.33.01 Driver Version: 440.33.01 CUDA Version:
10.2
* Failing:
TITAN V | NVIDIA-SMI 455.23.05 Driver Version: 455.23.05 CUDA Version:
11.1
Tesla V100-SXM2...| NVIDIA-SMI 410.79 Driver Version: 410.79 CUDA
Version: 10.0
Quadro GV100 | NVIDIA-SMI 460.32.03 Driver Version: 460.32.03 CUDA
Version: 11.2
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (5 preceding siblings ...)
2021-05-10 16:21 ` burnus at gcc dot gnu.org
@ 2021-05-10 17:37 ` vries at gcc dot gnu.org
2021-05-10 19:56 ` burnus at gcc dot gnu.org
` (8 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-10 17:37 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
CC| |tschwinge at gcc dot gnu.org
--- Comment #7 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #6)
> * Failing:
> TITAN V | NVIDIA-SMI 455.23.05 Driver Version: 455.23.05 CUDA Version:
> 11.1
Volta.
> Tesla V100-SXM2...| NVIDIA-SMI 410.79 Driver Version: 410.79 CUDA
> Version: 10.0
Volta.
> Quadro GV100 | NVIDIA-SMI 460.32.03 Driver Version: 460.32.03 CUDA
> Version: 11.2
and Volta.
I don't have a volta card, unfortunately.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (6 preceding siblings ...)
2021-05-10 17:37 ` vries at gcc dot gnu.org
@ 2021-05-10 19:56 ` burnus at gcc dot gnu.org
2021-05-11 9:32 ` vries at gcc dot gnu.org
` (7 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-05-10 19:56 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #8 from Tobias Burnus <burnus at gcc dot gnu.org> ---
I am wondering whether it has something to do with shfl now requiring .sync,
especially for sm_70. (Non-sync version was deprecated in ISA 6.0 and for sm_70
removed in ISA 6.4.)
Cross ref: PR96005.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (7 preceding siblings ...)
2021-05-10 19:56 ` burnus at gcc dot gnu.org
@ 2021-05-11 9:32 ` vries at gcc dot gnu.org
2021-05-11 12:58 ` burnus at gcc dot gnu.org
` (6 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-11 9:32 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #9 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #8)
> I am wondering whether it has something to do with shfl now requiring .sync,
> especially for sm_70. (Non-sync version was deprecated in ISA 6.0 and for
> sm_70 removed in ISA 6.4.)
> Cross ref: PR96005.
Could be. Posted a tentative patch at that PR, you could try that out.
Still, only root cause analysis can get us ... the root cause. I can try to do
that if you provide me with access to one of the volta machines. Otherwise,
can you or somebody else with access to these machines do that? I could coach
the effort a bit, if necessary.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (8 preceding siblings ...)
2021-05-11 9:32 ` vries at gcc dot gnu.org
@ 2021-05-11 12:58 ` burnus at gcc dot gnu.org
2021-05-11 13:00 ` jakub at gcc dot gnu.org
` (5 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-05-11 12:58 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
Tobias Burnus <burnus at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
CC| |jakub at gcc dot gnu.org
--- Comment #10 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Oddly, if one replaces
#pragma omp target teams distribute parallel for reduction(||: orfc) map(orfc)
by
#pragma omp target map(orfc)
{
#pragma teams distribute parallel for reduction(||: orfc)
for (int i=0; i < N; ++i)
orfc = orfc || rf[i];
}
the generated code vastly differs by not generating any 'omp teams distribute
for' code but just the plain for loop, which works.
Shouldn't this both generate the same code? (The omp-lower dump shows 'omp
target num_teams(0)' for the combined statement and 'omp target num_teams(1)'
for the separate statement, whatever that clause means for a target directive.)
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (9 preceding siblings ...)
2021-05-11 12:58 ` burnus at gcc dot gnu.org
@ 2021-05-11 13:00 ` jakub at gcc dot gnu.org
2021-05-12 11:21 ` vries at gcc dot gnu.org
` (4 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: jakub at gcc dot gnu.org @ 2021-05-11 13:00 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #11 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #10)
> Oddly, if one replaces
>
> #pragma omp target teams distribute parallel for reduction(||: orfc)
> map(orfc)
>
> by
>
> #pragma omp target map(orfc)
> {
> #pragma teams distribute parallel for reduction(||: orfc)
Missing omp before teams.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (10 preceding siblings ...)
2021-05-11 13:00 ` jakub at gcc dot gnu.org
@ 2021-05-12 11:21 ` vries at gcc dot gnu.org
2021-05-17 15:50 ` vries at gcc dot gnu.org
` (3 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-12 11:21 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #12 from Tom de Vries <vries at gcc dot gnu.org> ---
After investigation by Tobias, this looks like an instance of PR96932.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (11 preceding siblings ...)
2021-05-12 11:21 ` vries at gcc dot gnu.org
@ 2021-05-17 15:50 ` vries at gcc dot gnu.org
2021-05-17 18:21 ` cvs-commit at gcc dot gnu.org
` (2 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-17 15:50 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #13 from Tom de Vries <vries at gcc dot gnu.org> ---
Posted: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570508.html
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (12 preceding siblings ...)
2021-05-17 15:50 ` vries at gcc dot gnu.org
@ 2021-05-17 18:21 ` cvs-commit at gcc dot gnu.org
2021-05-17 18:23 ` vries at gcc dot gnu.org
2021-05-17 18:24 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: cvs-commit at gcc dot gnu.org @ 2021-05-17 18:21 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #14 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:58f7c7e098b79c96403c8341823ec3ba1e8b3945
commit r12-846-g58f7c7e098b79c96403c8341823ec3ba1e8b3945
Author: Tom de Vries <tdevries@suse.de>
Date: Mon May 17 10:11:52 2021 +0200
[nvptx] Handle memmodel for atomic ops
The atomic ops in nvptx.md have memmodel arguments, which are currently
ignored.
Handle these, fixing test-case fails libgomp.c-c++-common/reduction-{5,6}.c
on volta.
Tested libgomp on x86_64-linux with nvptx accelerator.
gcc/ChangeLog:
2021-05-17 Tom de Vries <tdevries@suse.de>
PR target/100497
* config/nvptx/nvptx-protos.h (nvptx_output_atomic_insn): Declare
* config/nvptx/nvptx.c (nvptx_output_barrier)
(nvptx_output_atomic_insn): New function.
(nvptx_print_operand): Add support for 'B'.
* config/nvptx/nvptx.md: Use nvptx_output_atomic_insn for atomic
insns.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (13 preceding siblings ...)
2021-05-17 18:21 ` cvs-commit at gcc dot gnu.org
@ 2021-05-17 18:23 ` vries at gcc dot gnu.org
2021-05-17 18:24 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-17 18:23 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Status|UNCONFIRMED |RESOLVED
Resolution|--- |FIXED
Target Milestone|--- |12.0
--- Comment #15 from Tom de Vries <vries at gcc dot gnu.org> ---
Patch committed, marking-resolved-fixed.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/100497] [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems
2021-05-10 9:00 [Bug target/100497] New: [OpenMP][nvptx] libgomp.c-c++-common/reduction-5.c - fails on some nvptx systems burnus at gcc dot gnu.org
` (14 preceding siblings ...)
2021-05-17 18:23 ` vries at gcc dot gnu.org
@ 2021-05-17 18:24 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-17 18:24 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100497
--- Comment #16 from Tom de Vries <vries at gcc dot gnu.org> ---
*** Bug 96932 has been marked as a duplicate of this bug. ***
^ permalink raw reply [flat|nested] 17+ messages in thread