public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Tobias Burnus <burnus@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Disable SIMT for user-defined reduction Date: Fri, 14 May 2021 08:50:53 +0000 (GMT) [thread overview] Message-ID: <20210514085053.EA053385E447@sourceware.org> (raw) https://gcc.gnu.org/g:6723e6700a3178f7e0fe5696325b25ccb81e5d09 commit 6723e6700a3178f7e0fe5696325b25ccb81e5d09 Author: Tom de Vries <tdevries@suse.de> Date: Fri May 14 09:24:47 2021 +0200 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: --- gcc/ChangeLog.omp | 9 +++++++++ gcc/omp-low.c | 13 +++++++++++++ libgomp/ChangeLog.omp | 8 ++++++++ libgomp/testsuite/libgomp.c/target-44.c | 27 +++++++++++++++++++++++++++ 4 files changed, 57 insertions(+) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 407ecb406f5..8e306a70480 100644 --- 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 f0cc49c2048..5eb8aee9c5e 100644 --- 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 96fc5ead0a0..e0f1ec0b50c 100644 --- 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 index 00000000000..13e0c757845 --- /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; +}
reply other threads:[~2021-05-14 8:50 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20210514085053.EA053385E447@sourceware.org \ --to=burnus@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).