https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104783
Bug ID: 104783
Summary: [nvptx, openmp] Hang/abort with atomic update in simd
construct
Product: gcc
Version: 12.0
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
Target Milestone: ---
Minimized from
https://github.com/TApplencourt/OvO/blob/master/test_src/cpp/hierarchical_parallelism/atomic_add-float/target_teams_distribute__parallel_for__simd.cpp
.
Test-case:
...
$ cat libgomp/testsuite/libgomp.c/target_teams_distribute__parallel_for__simd.c
/* { dg-options "-O2" } */
int
main (void)
{
const unsigned expected_value = 1;
unsigned counter_N0 = 0;
#pragma omp target map(tofrom: counter_N0)
#pragma omp simd
for (int i = 0 ; i < 1 ; i++ )
{
#pragma omp atomic update
counter_N0 = counter_N0 + 1 ;
}
return 0;
}
...
With target board unix/-foffload=-mptx=3.1:
...
libgomp: cuCtxSynchronize error: unspecified launch failure (perhaps abort was
called)
libgomp: cuMemFree_v2 error: unspecified launch failure
libgomp: device finalization failed
FAIL: libgomp.c/target_teams_distribute__parallel_for__simd.c execution test
...
With target board unix: hangs.
With dg-options -O0:
...
PASS: libgomp.c/target_teams_distribute__parallel_for__simd.c execution test
...
