diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 7b122059c6e83033b7277e714ebe2e407390f401..1f14c4b1d6946cf486cdb979e3ed036cc3296178 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -4385,6 +4385,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/testsuite/libgomp.c/target-44.c b/libgomp/testsuite/libgomp.c/target-44.c new file mode 100644 index 0000000000000000000000000000000000000000..13e0c757845702372f8f1ac9eafd09f02936139d --- /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; +}