[openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR
Consider test-case pr104952-1.c, included in this commit, containing: ... #pragma omp target map(tofrom:result) map(to:arr) #pragma omp simd reduction(||: result) ... When run on x86_64 with nvptx accelerator, the test-case either aborts or hangs. The reduction clause is translated by the SIMT code (active for nvptx) as a butterfly reduction loop with this butterfly shuffle / update pair: ... D.2163 = D.2163 || .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164) ... in the loop body. The problem is that the butterfly shuffle is possibly not executed, while it needs to be executed unconditionally. Fix this by translating instead as: ... D.tmp_bfly = .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164) D.2163 = D.2163 || D.tmp_bfly ... Tested on x86_64-linux with nvptx accelerator. gcc/ChangeLog: 2022-03-17 Tom de Vries <tdevries@suse.de> PR target/104952 * omp-low.cc (lower_rec_input_clauses): Make sure GOMP_SIMT_XCHG_BFLY is executed unconditionally. libgomp/ChangeLog: 2022-03-17 Tom de Vries <tdevries@suse.de> PR target/104952 * testsuite/libgomp.c/pr104952-1.c: New test. * testsuite/libgomp.c/pr104952-2.c: New test.
This commit is contained in:
parent
6393122d27
commit
093cdadbce
|
@ -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;
|
||||
|
|
|
@ -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;
|
||||
}
|
|
@ -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;
|
||||
}
|
Loading…
Reference in New Issue