* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
@ 2020-09-22 8:10 ` vries at gcc dot gnu.org
2021-05-11 9:23 ` vries at gcc dot gnu.org
` (4 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2020-09-22 8:10 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #0)
> Currently, we're at ptx isa v3.1:
> ...
> static void
> nvptx_file_start (void)
> {
> fputs ("// BEGIN PREAMBLE\n", asm_out_file);
> fputs ("\t.version\t3.1\n", asm_out_file);
> ...
>
> Using a newer isa could give some benefits, f.i. starting PTX ISA 6.3 we
> have atom.cas.b16.
As well as .alias, see PR97102.
A standalone test run with the patch from PR97102 comment 4 which bumps ISA to
6.3, runs into a few problems.
In libgomp testing: shfl deprecated, should use shfl.sync.
^ permalink raw reply [flat|nested] 7+ messages in thread* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
2020-09-22 8:10 ` [Bug target/96005] " vries at gcc dot gnu.org
@ 2021-05-11 9:23 ` vries at gcc dot gnu.org
2021-05-12 6:12 ` vries at gcc dot gnu.org
` (3 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-11 9:23 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
On my usual machine, using system cuda I don't get beyond 6.1:
...
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 7a7a9130e84..ecf3803df3c 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -5309,7 +5309,9 @@ static void
nvptx_file_start (void)
{
fputs ("// BEGIN PREAMBLE\n", asm_out_file);
- fputs ("\t.version\t3.1\n", asm_out_file);
+ /* PTX ISA version 6.1 is the current version for ptxas from cuda 9.1,
+ which is the current system release on ubuntu 18.04.5. */
+ fputs ("\t.version\t6.1\n", asm_out_file);
if (TARGET_SM35)
fputs ("\t.target\tsm_35\n", asm_out_file);
else
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 0f15609ee4b..39a695d275c 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -1452,7 +1452,7 @@
(match_operand:SI 3 "const_int_operand" "n")]
UNSPEC_SHUFFLE))]
""
- "%.\\tshfl%S3.b32\\t%0, %1, %2, 31;")
+ "%.\\tshfl.sync%S3.b32\\t%0, %1, %2, 31, 0xffffffff;")
(define_insn "nvptx_vote_ballot"
[(set (match_operand:SI 0 "nvptx_register_operand" "=R")
...
^ permalink raw reply [flat|nested] 7+ messages in thread* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
2020-09-22 8:10 ` [Bug target/96005] " vries at gcc dot gnu.org
2021-05-11 9:23 ` vries at gcc dot gnu.org
@ 2021-05-12 6:12 ` vries at gcc dot gnu.org
2021-05-12 10:59 ` vries at gcc dot gnu.org
` (2 subsequent siblings)
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-12 6:12 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #2)
> On my usual machine, using system cuda I don't get beyond 6.1:
Upgraded to ubuntu 20.4, giving me system cuda 10.1, which allows me to use isa
6.3. Now testing (using downgrade of driver from 460 -> 450 due to PR99932).
^ permalink raw reply [flat|nested] 7+ messages in thread* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
` (2 preceding siblings ...)
2021-05-12 6:12 ` vries at gcc dot gnu.org
@ 2021-05-12 10:59 ` vries at gcc dot gnu.org
2021-05-12 16:38 ` cvs-commit at gcc dot gnu.org
2021-05-12 16:39 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-12 10:59 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 50800
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=50800&action=edit
Tentative patch
^ permalink raw reply [flat|nested] 7+ messages in thread* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
` (3 preceding siblings ...)
2021-05-12 10:59 ` vries at gcc dot gnu.org
@ 2021-05-12 16:38 ` cvs-commit at gcc dot gnu.org
2021-05-12 16:39 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: cvs-commit at gcc dot gnu.org @ 2021-05-12 16:38 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
--- Comment #5 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:
https://gcc.gnu.org/g:2a1586401a21dcd43e0f904bb6eec26c8b2f366b
commit r12-751-g2a1586401a21dcd43e0f904bb6eec26c8b2f366b
Author: Tom de Vries <tdevries@suse.de>
Date: Wed May 12 12:40:37 2021 +0200
[nvptx] Add -mptx=3.1/6.3
Add nvptx option -mptx that sets the ptx ISA version. This is currently
hardcoded to 3.1.
Tested libgomp on x86_64-linux with nvptx accelerator, both with default
set to
3.1 and 6.3.
gcc/ChangeLog:
2021-05-12 Tom de Vries <tdevries@suse.de>
PR target/96005
* config/nvptx/nvptx-opts.h (enum ptx_version): New enum.
* config/nvptx/nvptx.c (nvptx_file_start): Print .version according
to ptx_version_option.
* config/nvptx/nvptx.h (TARGET_PTX_6_3): Define.
* config/nvptx/nvptx.md (define_insn "nvptx_shuffle<mode>")
(define_insn "nvptx_vote_ballot"): Use sync variant for
TARGET_PTX_6_3.
* config/nvptx/nvptx.opt (ptx_version): Add enum.
(mptx): Add option.
* doc/invoke.texi (Nvidia PTX Options): Add mptx item.
^ permalink raw reply [flat|nested] 7+ messages in thread* [Bug target/96005] Add possibility to use newer ptx isa
2020-06-30 14:36 [Bug target/96005] New: Add possibility to use newer ptx isa vries at gcc dot gnu.org
` (4 preceding siblings ...)
2021-05-12 16:38 ` cvs-commit at gcc dot gnu.org
@ 2021-05-12 16:39 ` vries at gcc dot gnu.org
5 siblings, 0 replies; 7+ messages in thread
From: vries at gcc dot gnu.org @ 2021-05-12 16:39 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Status|UNCONFIRMED |RESOLVED
Resolution|--- |FIXED
Target Milestone|--- |12.0
--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Patch committed, marking resolved-fixed.
^ permalink raw reply [flat|nested] 7+ messages in thread