From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 48) id 210D3385802D; Thu, 15 Apr 2021 08:02:07 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 210D3385802D From: "vries at gcc dot 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 X-Bugzilla-Reason: CC X-Bugzilla-Type: changed X-Bugzilla-Watch-Reason: None X-Bugzilla-Product: gcc X-Bugzilla-Component: target X-Bugzilla-Version: 11.0 X-Bugzilla-Keywords: openmp, wrong-code X-Bugzilla-Severity: normal X-Bugzilla-Who: vries at gcc dot gnu.org X-Bugzilla-Status: UNCONFIRMED X-Bugzilla-Resolution: X-Bugzilla-Priority: P3 X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org X-Bugzilla-Target-Milestone: --- X-Bugzilla-Flags: X-Bugzilla-Changed-Fields: Message-ID: In-Reply-To: References: Content-Type: text/plain; charset="UTF-8" Content-Transfer-Encoding: quoted-printable X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/ Auto-Submitted: auto-generated MIME-Version: 1.0 X-BeenThere: gcc-bugs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-bugs mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 15 Apr 2021 08:02:07 -0000 https://gcc.gnu.org/bugzilla/show_bug.cgi?id=3D99555 --- Comment #4 from Tom de Vries --- 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=20 * 0 - 0 1 Active 0x00000010 (1,1,1) (32,8,1) main$_omp_fn()= =20 ... 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 Filenam= e=20 Line=20 Kernel 0 * (0,0,0) (0,0,0) (0,0,0) (0,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (0,1,0) (0,0,0) (0,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (1,0,0) (0,0,0) (1,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (1,1,0) (0,0,0) (1,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (2,0,0) (0,0,0) (2,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (2,1,0) (0,0,0) (2,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (3,0,0) (0,0,0) (3,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (3,1,0) (0,0,0) (3,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (4,0,0) (0,0,0) (4,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (4,1,0) (0,0,0) (4,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (5,0,0) (0,0,0) (5,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (5,1,0) (0,0,0) (5,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (6,0,0) (0,0,0) (6,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (6,1,0) (0,0,0) (6,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (7,0,0) (0,0,0) (7,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (7,1,0) (0,0,0) (7,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (8,0,0) (0,0,0) (8,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (8,1,0) (0,0,0) (8,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (9,0,0) (0,0,0) (9,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (9,1,0) (0,0,0) (9,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (10,0,0) (0,0,0) (10,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (10,1,0) (0,0,0) (10,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (11,0,0) (0,0,0) (11,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (11,1,0) (0,0,0) (11,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (12,0,0) (0,0,0) (12,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (12,1,0) (0,0,0) (12,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (13,0,0) (0,0,0) (13,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (13,1,0) (0,0,0) (13,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (14,0,0) (0,0,0) (14,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (14,1,0) (0,0,0) (14,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (15,0,0) (0,0,0) (15,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (15,1,0) (0,0,0) (15,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (16,0,0) (0,0,0) (16,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (16,1,0) (0,0,0) (16,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (17,0,0) (0,0,0) (17,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (17,1,0) (0,0,0) (17,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (18,0,0) (0,0,0) (18,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (18,1,0) (0,0,0) (18,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (19,0,0) (0,0,0) (19,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (19,1,0) (0,0,0) (19,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (20,0,0) (0,0,0) (20,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (20,1,0) (0,0,0) (20,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (21,0,0) (0,0,0) (21,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (21,1,0) (0,0,0) (21,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (22,0,0) (0,0,0) (22,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (22,1,0) (0,0,0) (22,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (23,0,0) (0,0,0) (23,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (23,1,0) (0,0,0) (23,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (24,0,0) (0,0,0) (24,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (24,1,0) (0,0,0) (24,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (25,0,0) (0,0,0) (25,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (25,1,0) (0,0,0) (25,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (26,0,0) (0,0,0) (26,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (26,1,0) (0,0,0) (26,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (27,0,0) (0,0,0) (27,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (27,1,0) (0,0,0) (27,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (28,0,0) (0,0,0) (28,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (28,1,0) (0,0,0) (28,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (29,0,0) (0,0,0) (29,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (29,1,0) (0,0,0) (29,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (30,0,0) (0,0,0) (30,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (30,1,0) (0,0,0) (30,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 (0,0,0) (31,0,0) (0,0,0) (31,0,0) 1 0x0000000000b5f638 n/= a=20=20=20 0=20 (0,0,0) (31,1,0) (0,0,0) (31,7,0) 7 0x0000000000b2f350 n/= a=20=20=20 0=20 ... 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 =3D> 0x0000000000b5f638 <+1656>: MEMBAR.CTA ... and: ... 0x0000000000b2f340 <+4000>: 0x0000000000b2f348 <+4008>: BAR.SYNC 0x0, R4 =3D> 0x0000000000b2f350 <+4016>: MEMBAR.CTA ... Printing the registers for the thread amount operand gives us: ... (cuda-gdb) p $R0 $1 =3D 256 ... and: ... (cuda-gdb) cuda thread (0,1,0) ... (cuda-gdb) p $R4 $2 =3D 256 ... So we seem to be stuck at two barrier instructions both requiring 256 threa= ds, but a different logical barrier.=