[Bug target/96005] New: Add possibility to use newer ptx isa

[Bug target/96005] New: Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005

            Bug ID: 96005
           Summary: Add possibility to use newer ptx isa
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Severity: enhancement
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

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.

[Bug target/96005] Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

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.

[Bug target/96005] Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

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")
...

[Bug target/96005] Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

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).

[Bug target/96005] Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

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

[Bug target/96005] Add possibility to use newer ptx isa

[cvs-commit at gcc dot gnu.org]

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.

[Bug target/96005] Add possibility to use newer ptx isa

[vries at gcc dot gnu.org]

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.