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.
