https://gcc.gnu.org/bugzilla/show_bug.cgi?id=128414
Bug ID: 128414
Summary: [17 Regression][nvptx] wrong code with 'target teams'
running too often and failing with cuCtxSynchronize
Product: gcc
Version: 17.0
Status: UNCONFIRMED
Keywords: wrong-code
Severity: normal
Priority: P3
Component: libgomp
Assignee: unassigned at gcc dot gnu.org
Reporter: burnus at gcc dot gnu.org
CC: jakub at gcc dot gnu.org, matmal01 at gcc dot gnu.org,
tschwinge at gcc dot gnu.org
Target Milestone: ---
Target: nvptx
The following code is supposed to offload to the GPU and print there
0 - 0 - 0
and exit with exit status 0. That happens with nvptx offloading an my
GCC 16 distro compiler.
The code is:
#include <omp.h>
int main () {
#pragma omp target teams num_teams(1)
{
#pragma omp parallel if(0)
{
__builtin_printf ("%d - %d - %d\n",
omp_get_device_num (), omp_get_team_num(), omp_get_thread_num());
}
}
}
However, it now does with nvptx offloading:
0 - 0 - 0
0 - 0 - 0
0 - 0 - 0
0 - 0 - 0
(in total 31 such lines)
libgomp: cuCtxSynchronize error: unspecified launch failure (perhaps abort
was called)
* * *
That's with r17-5896-g08f84a62a4cf14 (i.e. fairly recent with Matthew's patches
applied).
I tried to revert Matthew's five patches as the first
72d86fead1f libgomp: Inline `gomp_barrier_state_is_incremented`
libgomp/config/nvptx/bar.h | 13 ++-----------
and the second
5ffc8847dfc libgomp: New barriers in new linux/futex_waitv target
libgomp/config/nvptx/bar.c | 36 +-
libgomp/config/nvptx/bar.h | 63 +-
libgomp/config/nvptx/team.c | 3 +-
touch nvptx code -- especially as they look as if they might have caused the
issue.
However, reverting all five patches and re-compiling did not fix it.
Hence, either the regression is older - or I messed up with my reverting -
or my tree is broken for some reasons [i.e. needs a clean bootstrap and proper
regression hunting] - but I don't feel like regression hunting this evening.
