From: Cesar Philippidis <cesar@codesourcery.com>
To: Tom de Vries <Tom_deVries@mentor.com>
Cc: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>,
Thomas Schwinge <thomas@codesourcery.com>
Subject: Re: [og7] Update nvptx_fork/join barrier placement
Date: Mon, 19 Mar 2018 18:09:00 -0000 [thread overview]
Message-ID: <99134339-58fa-0134-db6c-06a78ea376b7@codesourcery.com> (raw)
In-Reply-To: <c6f53eb4-dd5d-a539-907a-212dc0d1cd6c@mentor.com>
[-- Attachment #1: Type: text/plain, Size: 862 bytes --]
On 03/19/2018 10:02 AM, Tom de Vries wrote:
> On 03/19/2018 03:55 PM, Cesar Philippidis wrote:
>>> Note that this changes ordering of the vector-neutering jump and
>>> worker-neutering jump at the end. In principle, this should not be
>>> harmful, but it violates the invariant that vector-neutering
>>> branch-around code should be as short-lived as possible. So, this needs
>>> to be fixed.
>>>
>>> I've found this issue by adding verification of the neutering, as
>>> attached below.
>> ACK, thanks. I'll take a closer look at this.
>
> I've got a tentative patch at
> https://gcc.gnu.org/bugzilla/attachment.cgi?id=43707 ( PR84952 -
> "[nvptx] bar.sync generated in divergent code" ).
I attached my WIP patch. But, given that you've spent a lot of time on
this, I'll let you continue working on it. Just remember to backport any
fix to og7.
Thanks,
Cesar
[-- Attachment #2: trunk-barriers-new.diff --]
[-- Type: text/x-patch, Size: 2341 bytes --]
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index a6f444340fd..0d288cb81ba 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4037,6 +4037,22 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
return;
}
+ /* NVPTX_BARSYNC barriers are placed immediately before NVPTX_JOIN
+ in order to ensure that all of the threads in a CTA reach the
+ barrier. Don't nueter BLOCK if head is NVPTX_BARSYNC and tail is
+ NVPTX_JOIN. */
+ if (from == to
+ && recog_memoized (head) == CODE_FOR_nvptx_barsync
+ && recog_memoized (tail) == CODE_FOR_nvptx_join)
+ return;
+
+ /* Adjust HEAD to point to the NVPTX_JOIN instruction after a
+ NVPTX_BARSYNC, so that any successive state neutering code does
+ not get placed before the dummy JOIN comment. */
+ if (recog_memoized (head) == CODE_FOR_nvptx_barsync
+ && recog_memoized (NEXT_INSN (head)) == CODE_FOR_nvptx_join)
+ head = NEXT_INSN (head);
+
/* Insert the vector test inside the worker test. */
unsigned mode;
rtx_insn *before = tail;
@@ -4057,7 +4073,23 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
br = gen_br_true (pred, label);
else
br = gen_br_true_uni (pred, label);
- emit_insn_before (br, head);
+
+ if (recog_memoized (head) == CODE_FOR_nvptx_forked
+ && recog_memoized (NEXT_INSN (head)) == CODE_FOR_nvptx_barsync)
+ {
+ head = NEXT_INSN (head);
+ emit_insn_after (br, head);
+ }
+ else if (recog_memoized (head) == CODE_FOR_nvptx_join)
+ {
+ if (recog_memoized (NEXT_INSN (head)) == CODE_FOR_br_true_uni
+ && mode == GOMP_DIM_VECTOR)
+ emit_insn_after (br, NEXT_INSN (head));
+ else
+ emit_insn_after (br, head);
+ }
+ else
+ emit_insn_before (br, head);
LABEL_NUSES (label)++;
if (tail_branch)
@@ -4276,7 +4308,7 @@ nvptx_process_pars (parallel *par)
nvptx_wpropagate (true, par->forked_block, par->fork_insn);
/* Insert begin and end synchronizations. */
emit_insn_after (nvptx_wsync (false), par->forked_insn);
- emit_insn_before (nvptx_wsync (true), par->joining_insn);
+ emit_insn_before (nvptx_wsync (true), par->join_insn);
}
else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR))
nvptx_vpropagate (par->forked_block, par->forked_insn);
next prev parent reply other threads:[~2018-03-19 17:55 UTC|newest]
Thread overview: 9+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-03-08 23:31 Cesar Philippidis
2018-03-09 16:22 ` Tom de Vries
2018-03-09 16:55 ` Cesar Philippidis
2018-03-19 14:30 ` Tom de Vries
2018-03-19 15:24 ` Cesar Philippidis
2018-03-19 15:44 ` Tom de Vries
2018-03-19 17:04 ` Tom de Vries
2018-03-19 18:09 ` Cesar Philippidis [this message]
2018-03-20 10:41 ` [nvptx, PR84952, committed] Fix bar.sync position Tom de Vries
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=99134339-58fa-0134-db6c-06a78ea376b7@codesourcery.com \
--to=cesar@codesourcery.com \
--cc=Tom_deVries@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=thomas@codesourcery.com \
/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: link
Be 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).