[committed,openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR

Message ID 20220318144716.GA30949@delia.home
State Committed
Commit 093cdadbce30ce2d36846a05d979b8afc2eff618
Headers
Series [committed,openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR |

Commit Message

Tom de Vries March 18, 2022, 2:47 p.m. UTC
  Hi,

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.

Committed to trunk.

Thanks,
- Tom

[openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR

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.

---
 gcc/omp-low.cc                           |  5 ++++-
 libgomp/testsuite/libgomp.c/pr104952-1.c | 24 ++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c/pr104952-2.c | 22 ++++++++++++++++++++++
 3 files changed, 50 insertions(+), 1 deletion(-)
  

Patch

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index cfc63d6a104..392bb18bc5d 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 00000000000..a3bfb1e77df
--- /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 00000000000..7ab4bcdb8af
--- /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;
+}