Disable SIMT for user-defined reduction

The test-case included in this patch contains this target region:
...
  for (int i0 = 0 ; i0 < N0 ; i0++ )
    counter_N0.i += 1;
...

When running with nvptx accelerator, the counter variable is expected to
be N0 after the region, but instead is N0 / 32.  The problem is that rather
than getting the result for all warp lanes, we get it for just one lane.

This is caused by the implementation of SIMT being incomplete.  It handles
regular reductions, but appearantly not user-defined reductions.

For now, handle this by disabling SIMT in this case, specifically by setting
sctx->max_vf to 1.

Tested libgomp on x86_64-linux with nvptx accelerator.

gcc/ChangeLog:

2021-05-03  Tom de Vries  <tdevries@suse.de>

	PR target/100321
	* omp-low.c (lower_rec_input_clauses): Disable SIMT for user-defined
	reduction.

libgomp/ChangeLog:

2021-05-03  Tom de Vries  <tdevries@suse.de>

	PR target/100321
	* testsuite/libgomp.c/target-44.c: New test.

(cherry picked from commit f87990a2a8fc9e20d30462a0a4c9047582af0cd9)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 407ecb406f5185c0778ac516868ee383aedde389..8e306a704801748a62237880c0efa8b5e149029d 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2021-05-14  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backported from master:
+       2021-05-03  Tom de Vries  <tdevries@suse.de>
+
+       PR target/100321
+       * omp-low.c (lower_rec_input_clauses): Disable SIMT for user-defined
+       reduction.
+
 2021-05-14  Tobias Burnus  <tobias@codesourcery.com>
 
        Backported from master:

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f0cc49c2048e31aea9b0b9aa8e508b7ff19e580f..5eb8aee9c5e68c86d338619e512563642fa2860b 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -4710,6 +4710,19 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx,
                sctx->max_vf = lower_bound (sctx->max_vf, safe_len);
            }
        }
+      if (sctx->is_simt && !known_eq (sctx->max_vf, 1U))
+       {
+         for (tree c = gimple_omp_for_clauses (ctx->stmt); c;
+              c = OMP_CLAUSE_CHAIN (c))
+           if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
+               && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
+             {
+               /* UDR reductions are not supported yet for SIMT, disable
+                  SIMT.  */
+               sctx->max_vf = 1;
+               break;
+             }
+       }
       if (maybe_gt (sctx->max_vf, 1U))
        {
          sctx->idx = create_tmp_var (unsigned_type_node);

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 96fc5ead0a0913248f2835ff99c5f9e778342ee2..e0f1ec0b50c2874c3deb8e15c7dcfce766d89826 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2021-05-14  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backported from master:
+       2021-05-03  Tom de Vries  <tdevries@suse.de>
+
+       PR target/100321
+       * testsuite/libgomp.c/target-44.c: New test.
+
 2021-05-14  Tobias Burnus  <tobias@codesourcery.com>
 
        Backported from master:

diff --git a/libgomp/testsuite/libgomp.c/target-44.c b/libgomp/testsuite/libgomp.c/target-44.c
new file mode 100644 (file)
index 0000000..13e0c75
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-44.c
@@ -0,0 +1,27 @@
+/* { dg-additional-options "-foffload=-latomic" { target { offload_target_nvptx } } } */
+
+#include <stdlib.h>
+
+struct s
+{
+  int i;
+};
+
+#pragma omp declare reduction(+: struct s: omp_out.i += omp_in.i)
+
+int
+main (void)
+{
+  const int N0 = 32768;
+
+  struct s counter_N0 = { 0 };
+#pragma omp target
+#pragma omp for simd reduction(+: counter_N0)
+  for (int i0 = 0 ; i0 < N0 ; i0++ )
+    counter_N0.i += 1;
+
+  if (counter_N0.i != N0)
+    abort ();
+
+  return 0;
+}