KJ16609/gcc-reflection
1
0
Fork 0
mirror of https://forge.sourceware.org/marek/gcc.git synced 2026-02-22 12:00:11 -05:00
Code Issues Packages Projects Releases Wiki Activity

Fix gimple_copy for OpenMP atomic load/store [PR122281, PR105001]

Browse Source 
PR libgomp/122281
	PR middle-end/105001

gcc/ChangeLog:

	* gimple.cc (gimple_copy): Add missing unshare_expr for
	GIMPLE_OMP_ATOMIC_LOAD and GIMPLE_OMP_ATOMIC_STORE.
This commit is contained in:
Tobias Burnus
2025-11-03 12:21:30 +01:00
parent c10bd8ce54
commit e2cbcd1b27
2 changed files with 65 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

22
gcc/gimple.cc
View File

@@ -2283,6 +2283,28 @@ gimple_copy (gimple *stmt)
}
}
switch (gimple_code (stmt))
{
case GIMPLE_OMP_ATOMIC_LOAD:
{
gomp_atomic_load *g = as_a <gomp_atomic_load *> (copy);
gimple_omp_atomic_load_set_lhs (g,
unshare_expr (gimple_omp_atomic_load_lhs (g)));
gimple_omp_atomic_load_set_rhs (g,
unshare_expr (gimple_omp_atomic_load_rhs (g)));
break;
}
case GIMPLE_OMP_ATOMIC_STORE:
{
gomp_atomic_store *g = as_a <gomp_atomic_store *> (copy);
gimple_omp_atomic_store_set_val (g,
unshare_expr (gimple_omp_atomic_store_val (g)));
break;
}
default:
break;
}
/* Make copy of operands. */
for (i = 0; i < num_ops; i++)
gimple_set_op (copy, i, unshare_expr (gimple_op (stmt, i)));

43
libgomp/testsuite/libgomp.c/pr122281.c Normal file
View File

@@ -0,0 +1,43 @@
/* { dg-do run } */
/* { dg-additional-options "-O3" } */
/* PR libgomp/122281 */
/* PR middle-end/105001 */
/* If SIMT is supported, the inner 'omp simd' is duplicated into
one SIMT and one SIMD variant. SIMT is currently only supported
with nvidia GPUs. (This only happens with -O1 or higher.)
The duplication failed for the SIMD case as a tree was shared and
the initialization only happened in the SIMT branch, i.e. when
compiling for a SIMT-device, all non-SIMD (offload or host devices)
accesses failed (segfault) for the atomic update. */
#include <omp.h>
int __attribute__((noinline, noclone))
f(int *A, int n, int dev) {
int cnt = 0;
#pragma omp target map(cnt) device(dev)
{
#pragma omp parallel for simd
for (int i = 0; i < n; i++)
if (A[i] != 0)
{
#pragma omp atomic
cnt++;
}
}
return cnt;
}
int main() {
int n = 10;
int A[10] = {11,22,33,44,55,66,77,88,99,110};
/* Run over all devices, including the host; the host should be SIMD,
some non-host devices might be SIMT. */
for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
if (f (A, n, dev) != 10)
__builtin_abort();
}