On 5/6/21 12:30 PM, Jakub Jelinek wrote: > On Thu, May 06, 2021 at 12:17:03PM +0200, Tobias Burnus wrote: >> OpenMP: Fix SIMT for complex/float reduction with && and || >> >> gcc/ChangeLog: >> >> * omp-low.c (lower_rec_input_clauses): Also handle SIMT part >> for complex/float recution with && and ||. >> >> libgomp/ChangeLog: >> >> * testsuite/libgomp.c-c++-common/reduction-5.c: New test, testing >> complex/floating-point || + && recduction with 'omp target'. > > As the float/complex ||/&& reductions are IMHO just conformance issues, not > something anyone would actually use in meaningful code - floats or complex > aren't the most obvious or efficient holders of boolean values - I think > punting SIMT on those isn't a workaround, but the right solution. >
Ack. WIP patch below tries that approach and fixes the ICE, but this simple example still doesn't work: ... int main () { float andf = 1; #pragma omp target parallel reduction(&&: andf) for (int i=0; i < 1024; ++i) andf = andf && 0.0; if ((int)andf != 0) __builtin_abort (); return 0; } ... Thanks, - Tom
diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 26ceaf74b2d..d8f2487054f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -4389,14 +4389,28 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, { 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 (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION) + continue; + + if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) + { + /* UDR reductions are not supported yet for SIMT, disable + SIMT. */ + sctx->max_vf = 1; + break; + } + + if (truth_value_p (OMP_CLAUSE_REDUCTION_CODE (c)) + && TREE_CODE (TREE_TYPE (new_var)) != BOOLEAN_TYPE) + { + /* Doing boolean operations on non-boolean types is + for conformance only, it's not worth supporting this + for SIMT. */ + sctx->max_vf = 1; + break; + } + } } if (maybe_gt (sctx->max_vf, 1U)) {