From: Tom de Vries <tdevries@suse.de>
To: "Schwinge, Thomas" <Thomas_Schwinge@mentor.com>
Cc: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>
Subject: [nvptx, committed] Generalize bar.sync instruction
Date: Wed, 19 Dec 2018 10:33:00 -0000 [thread overview]
Message-ID: <34ce9b0d-899d-186f-93bc-c264ae309817@suse.de> (raw)
In-Reply-To: <1394d89c-896e-f6a3-5f9a-78e98b16e85c@suse.de>
[-- Attachment #1: Type: text/plain, Size: 193 bytes --]
[ was: Re: [nvptx] vector length patch series ]
On 14-12-18 20:58, Tom de Vries wrote:
> 0011-nvptx-Add-thread-count-parm-to-bar.sync.patch
Factored out this patch, committed.
Thanks,
- Tom
[-- Attachment #2: 0003-nvptx-Generalize-bar.sync-instruction.patch --]
[-- Type: text/x-patch, Size: 1670 bytes --]
[nvptx] Generalize bar.sync instruction
Allow the logical barrier operand of nvptx_barsync to be a register, and add a
thread count operand.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-12-17 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.md (nvptx_barsync): Add and handle operand.
* config/nvptx/nvptx.c (nvptx_wsync): Update call to gen_nvptx_barsync.
---
gcc/config/nvptx/nvptx.c | 2 +-
gcc/config/nvptx/nvptx.md | 10 ++++++++--
2 files changed, 9 insertions(+), 3 deletions(-)
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index a354811194c..1ad3ba92caa 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -3974,7 +3974,7 @@ nvptx_wpropagate (bool pre_p, bool is_call, basic_block block, rtx_insn *insn)
static rtx
nvptx_wsync (bool after)
{
- return gen_nvptx_barsync (GEN_INT (after));
+ return gen_nvptx_barsync (GEN_INT (after), GEN_INT (0));
}
#if WORKAROUND_PTXJIT_BUG
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index ca00b1d8073..f1f6fe0c404 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -1454,10 +1454,16 @@
[(set_attr "atomic" "true")])
(define_insn "nvptx_barsync"
- [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
+ [(unspec_volatile [(match_operand:SI 0 "nvptx_nonmemory_operand" "Ri")
+ (match_operand:SI 1 "const_int_operand")]
UNSPECV_BARSYNC)]
""
- "\\tbar.sync\\t%0;"
+ {
+ if (INTVAL (operands[1]) == 0)
+ return "\\tbar.sync\\t%0;";
+ else
+ return "\\tbar.sync\\t%0, %1;";
+ }
[(set_attr "predicable" "false")])
(define_expand "memory_barrier"
next prev parent reply other threads:[~2018-12-19 10:33 UTC|newest]
Thread overview: 32+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-04-23 11:24 [og7, nvptx, PR85486, committed] Force vl32 if calling vector-partitionable routines Tom de Vries
2018-09-18 20:22 ` [nvptx] vector length patch series Cesar Philippidis
2018-10-06 7:49 ` Tom de Vries
2018-10-29 20:09 ` Cesar Philippidis
2018-12-14 19:58 ` Tom de Vries
2018-12-17 21:29 ` [nvptx, committed] Rewrite nvptx_goacc_validate_dims to use predicate vars Tom de Vries
2018-12-17 21:46 ` [nvptx, committed] Unify C/Fortran routine handling in nvptx_goacc_validate_dims Tom de Vries
2019-02-22 11:23 ` Thomas Schwinge
2018-12-17 21:48 ` [nvptx] vector length patch series Tom de Vries
2018-12-17 21:51 ` [nvptx, committed] Add PTX_WARP_SIZE Tom de Vries
2018-12-17 21:52 ` [nvptx, committed] Move macro defs to top of nvptx.c Tom de Vries
2018-12-17 23:52 ` [nvptx] vector length patch series -- openacc parts Tom de Vries
2019-01-03 9:13 ` Tom de Vries
2018-12-19 10:28 ` [nvptx, committed] Use TARGET_SET_CURRENT_FUNCTION Tom de Vries
2018-12-19 10:31 ` [nvptx, committed] Only use one logical barrier resource Tom de Vries
2018-12-19 10:33 ` Tom de Vries [this message]
2018-12-19 10:35 ` [nvptx, committed] Rename worker_bcast variables to oacc_bcast Tom de Vries
2018-12-19 10:37 ` [nvptx, committed] Make nvptx state propagation function names more generic Tom de Vries
2018-12-19 10:38 ` [nvptx, committed] Use MAX, MIN, ROUND_UP macros Tom de Vries
2018-12-19 17:19 ` [nvptx, committed] Add PTX_CTA_SIZE Tom de Vries
2018-12-22 2:19 ` [nvptx] vector length patch series Tom de Vries
2019-01-05 11:19 ` Tom de Vries
2019-01-03 16:08 ` Tom de Vries
2019-01-03 16:20 ` Tom de Vries
2019-01-03 16:29 ` Tom de Vries
2019-01-07 8:52 ` [nvptx, committed] Fix libgomp.oacc-c-c++-common/vector-length-128-3.c Tom de Vries
2019-01-07 8:59 ` [nvptx, committed] Add support for a per-worker broadcast buffer and barrier Tom de Vries
2019-01-07 9:01 ` [nvptx] Don't emit barriers for empty loops -- fix Tom de Vries
2019-01-07 9:03 ` [nvptx] Handle large vector reductions Tom de Vries
2019-01-07 19:11 ` [nvptx, committed] Force vl32 if calling vector-partitionable routines Tom de Vries
2020-10-30 16:53 ` Thomas Schwinge
2019-01-08 23:00 ` [nvptx] vector length patch series 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=34ce9b0d-899d-186f-93bc-c264ae309817@suse.de \
--to=tdevries@suse.de \
--cc=Thomas_Schwinge@mentor.com \
--cc=gcc-patches@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: 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).