[Bug target/104952] [nvptx][OpenMP] wrong code with OR / AND reduction ('reduction(||:' and '&&') with SIMT

[vries at gcc dot gnu.org via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104952

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Keywords|                            |openmp

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
I can reproduce the problem.

I've made the simd explicit (I hope that's still valid openmp code):
...
$ cat libgomp/testsuite/libgomp.c/test.c
#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] = 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)
    return 1;

  return 0;
}
...

Easy workaround:
...
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d932d74cb03..bf6845d654e 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4641,6 +4641,15 @@ lower_rec_simd_input_clauses (tree new_var, omp_context
*ctx,
                  sctx->max_vf = 1;
                  break;
              }
+
+             if (OMP_CLAUSE_REDUCTION_CODE (c) == TRUTH_ANDIF_EXPR
+                 || OMP_CLAUSE_REDUCTION_CODE (c) == TRUTH_ORIF_EXPR)
+               {
+                 sctx->max_vf = 1;
+                 break;
+               }
            }
        }
       if (maybe_gt (sctx->max_vf, 1U))
...