diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index cfc63d6a104a..392bb18bc5dd 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -6743,7 +6743,10 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, x = build_call_expr_internal_loc (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_BFLY, TREE_TYPE (ivar), 2, ivar, simt_lane); - x = build2 (code, TREE_TYPE (ivar), ivar, x); + /* Make sure x is evaluated unconditionally. */ + tree bfly_var = create_tmp_var (TREE_TYPE (ivar)); + gimplify_assign (bfly_var, x, &llist[2]); + x = build2 (code, TREE_TYPE (ivar), ivar, bfly_var); gimplify_assign (ivar, x, &llist[2]); } tree ivar2 = ivar; diff --git a/libgomp/testsuite/libgomp.c/pr104952-1.c b/libgomp/testsuite/libgomp.c/pr104952-1.c new file mode 100644 index 000000000000..a3bfb1e77df2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr104952-1.c @@ -0,0 +1,24 @@ +#define N 32 + +static char arr[N]; + +int +main (void) +{ + unsigned int result = 0; + + for (unsigned int i = 0; i < N; ++i) + arr[i] = 0; + + arr[5] = 42; + +#pragma omp target map(tofrom:result) map(to:arr) +#pragma omp simd reduction(||: result) + for (unsigned int i = 0; i < N; ++i) + result = result || arr[i]; + + if (result != 1) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr104952-2.c b/libgomp/testsuite/libgomp.c/pr104952-2.c new file mode 100644 index 000000000000..7ab4bcdb8af3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr104952-2.c @@ -0,0 +1,22 @@ +#define N 32 + +static char arr[N]; + +int +main (void) +{ + unsigned int result = 2; + + for (unsigned int i = 0; i < N; ++i) + arr[i] = i + 1; + +#pragma omp target map(tofrom:result) map(to:arr) +#pragma omp simd reduction(&&: result) + for (unsigned int i = 0; i < N; ++i) + result = result && arr[i]; + + if (result != 1) + __builtin_abort (); + + return 0; +}