public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
From: "vries at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/99555] [OpenMP/nvptx] Execution-time hang for simple nested OpenMP 'target'/'parallel'/'task' constructs
Date: Thu, 15 Apr 2021 08:02:06 +0000 [thread overview]
Message-ID: <bug-99555-4-V1XUFFdKIa@http.gcc.gnu.org/bugzilla/> (raw)
In-Reply-To: <bug-99555-4@http.gcc.gnu.org/bugzilla/>
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99555
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
Investigated using cuda-gdb.
After typing ^c, we investigate the state:
...
(cuda-gdb) info cuda kernels
Kernel Parent Dev Grid Status SMs Mask GridDim BlockDim Invocation
* 0 - 0 1 Active 0x00000010 (1,1,1) (32,8,1) main$_omp_fn()
...
So, we have 256 threads in the CTA, or 8 warps.
The threads have the following state:
...
(cuda-gdb) info cuda threads
BlockIdx ThreadIdx To BlockIdx ThreadIdx Count Virtual PC Filename
Line
Kernel 0
* (0,0,0) (0,0,0) (0,0,0) (0,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (0,1,0) (0,0,0) (0,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (1,0,0) (0,0,0) (1,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (1,1,0) (0,0,0) (1,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (2,0,0) (0,0,0) (2,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (2,1,0) (0,0,0) (2,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (3,0,0) (0,0,0) (3,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (3,1,0) (0,0,0) (3,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (4,0,0) (0,0,0) (4,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (4,1,0) (0,0,0) (4,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (5,0,0) (0,0,0) (5,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (5,1,0) (0,0,0) (5,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (6,0,0) (0,0,0) (6,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (6,1,0) (0,0,0) (6,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (7,0,0) (0,0,0) (7,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (7,1,0) (0,0,0) (7,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (8,0,0) (0,0,0) (8,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (8,1,0) (0,0,0) (8,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (9,0,0) (0,0,0) (9,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (9,1,0) (0,0,0) (9,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (10,0,0) (0,0,0) (10,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (10,1,0) (0,0,0) (10,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (11,0,0) (0,0,0) (11,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (11,1,0) (0,0,0) (11,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (12,0,0) (0,0,0) (12,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (12,1,0) (0,0,0) (12,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (13,0,0) (0,0,0) (13,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (13,1,0) (0,0,0) (13,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (14,0,0) (0,0,0) (14,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (14,1,0) (0,0,0) (14,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (15,0,0) (0,0,0) (15,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (15,1,0) (0,0,0) (15,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (16,0,0) (0,0,0) (16,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (16,1,0) (0,0,0) (16,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (17,0,0) (0,0,0) (17,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (17,1,0) (0,0,0) (17,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (18,0,0) (0,0,0) (18,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (18,1,0) (0,0,0) (18,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (19,0,0) (0,0,0) (19,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (19,1,0) (0,0,0) (19,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (20,0,0) (0,0,0) (20,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (20,1,0) (0,0,0) (20,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (21,0,0) (0,0,0) (21,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (21,1,0) (0,0,0) (21,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (22,0,0) (0,0,0) (22,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (22,1,0) (0,0,0) (22,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (23,0,0) (0,0,0) (23,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (23,1,0) (0,0,0) (23,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (24,0,0) (0,0,0) (24,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (24,1,0) (0,0,0) (24,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (25,0,0) (0,0,0) (25,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (25,1,0) (0,0,0) (25,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (26,0,0) (0,0,0) (26,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (26,1,0) (0,0,0) (26,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (27,0,0) (0,0,0) (27,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (27,1,0) (0,0,0) (27,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (28,0,0) (0,0,0) (28,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (28,1,0) (0,0,0) (28,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (29,0,0) (0,0,0) (29,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (29,1,0) (0,0,0) (29,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (30,0,0) (0,0,0) (30,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (30,1,0) (0,0,0) (30,7,0) 7 0x0000000000b2f350 n/a
0
(0,0,0) (31,0,0) (0,0,0) (31,0,0) 1 0x0000000000b5f638 n/a
0
(0,0,0) (31,1,0) (0,0,0) (31,7,0) 7 0x0000000000b2f350 n/a
0
...
I seems that we're stuck at two locations, one warp in one and 7 warps in
another. Here (in thread 0,0,0):
...
(cuda-gdb) bt
#0 0x0000000000b5f638 in gomp_team_barrier_wait_end ()
#1 0x0000000000a9e638 in gomp_team_barrier_wait_final ()
#2 0x0000000000b31ad8 in gomp_team_end ()
#3 0x0000000000b394d8 in GOMP_parallel_end ()
#4 0x0000000000a7e620 in GOMP_parallel ()
#5 0x0000000000b48cc0 in main$_omp_fn$0$impl ()
#6 0x0000000000b2f020 in gomp_nvptx_main ()
#7 0x0000000000b4a2c0 in main$_omp_fn<<<(1,1,1),(32,8,1)>>> ()
...
and here:
...
(cuda-gdb) cuda thread (0,1,0)
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,1,0),
device 0, sm 4, warp 1, lane 0]
0x0000000000b2f350 in gomp_nvptx_main ()
(cuda-gdb) bt
#0 0x0000000000b2f350 in gomp_nvptx_main ()
#1 0x0000000000b4a2c0 in main$_omp_fn<<<(1,1,1),(32,8,1)>>> ()
...
Looking at the specific addresses, we have two bar.sync insns:
...
0x0000000000b5f630 <+1648>: BAR.SYNC 0x1, R0
=> 0x0000000000b5f638 <+1656>: MEMBAR.CTA
...
and:
...
0x0000000000b2f340 <+4000>:
0x0000000000b2f348 <+4008>: BAR.SYNC 0x0, R4
=> 0x0000000000b2f350 <+4016>: MEMBAR.CTA
...
Printing the registers for the thread amount operand gives us:
...
(cuda-gdb) p $R0
$1 = 256
...
and:
...
(cuda-gdb) cuda thread (0,1,0)
...
(cuda-gdb) p $R4
$2 = 256
...
So we seem to be stuck at two barrier instructions both requiring 256 threads,
but a different logical barrier.
next prev parent reply other threads:[~2021-04-15 8:02 UTC|newest]
Thread overview: 17+ messages / expand[flat|nested] mbox.gz Atom feed top
2021-03-11 16:36 [Bug target/99555] New: " tschwinge at gcc dot gnu.org
2021-03-12 15:53 ` [Bug target/99555] " vries at gcc dot gnu.org
2021-03-25 12:00 ` cvs-commit at gcc dot gnu.org
2021-03-29 8:41 ` cvs-commit at gcc dot gnu.org
2021-04-15 8:02 ` vries at gcc dot gnu.org [this message]
2021-04-15 9:14 ` cvs-commit at gcc dot gnu.org
2021-04-17 8:07 ` vries at gcc dot gnu.org
2021-04-19 10:44 ` vries at gcc dot gnu.org
2021-04-19 11:15 ` vries at gcc dot gnu.org
2021-04-19 15:39 ` vries at gcc dot gnu.org
2021-04-20 11:24 ` vries at gcc dot gnu.org
2022-02-22 14:53 ` cvs-commit at gcc dot gnu.org
2022-02-22 14:54 ` vries at gcc dot gnu.org
2022-03-17 12:16 ` tschwinge at gcc dot gnu.org
2022-05-13 13:16 ` tschwinge at gcc dot gnu.org
2022-09-06 13:32 ` vries at gcc dot gnu.org
2022-12-21 13:59 ` cvs-commit at gcc dot gnu.org
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=bug-99555-4-V1XUFFdKIa@http.gcc.gnu.org/bugzilla/ \
--to=gcc-bugzilla@gcc.gnu.org \
--cc=gcc-bugs@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).