[Bug target/97102] New: [nvptx] PTX JIT compilation failed when using aliases

[Bug target/97102] New: [nvptx] PTX JIT compilation failed when using aliases

[burnus at gcc dot gnu.org]

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

            Bug ID: 97102
           Summary: [nvptx] PTX JIT compilation failed when using aliases
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Keywords: openmp
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: vries at gcc dot gnu.org
  Target Milestone: ---

[Found when testing a patch for PR96390, which works with GCN offloading but
not with nvptx]

When compiling the following with -fopenmp
--------------------------------------
#pragma omp declare target
int foo () { return 42; }
int bar () __attribute__((alias ("foo")));
#pragma omp end declare target

int
main ()
{
  int n;
  #pragma omp target map(from:n)
    n = bar ();
  if (n != 42)
    __builtin_abort ();
}
--------------------------------------

it fails with when running with:

libgomp: Link error log ptxas application ptx input, line 89; error   : Label
expected for argument 0 of instruction 'call'
ptxas application ptx input, line 89; fatal   : Call target not recognized
ptxas fatal   : Ptx assembly aborted due to errors

libgomp: cuLinkAddData (ptx_code) error: a PTX JIT compilation failed


The generated PTX has:
------------
// BEGIN GLOBAL FUNCTION DECL: foo
.visible .func (.param.u32 %value_out) foo;

// BEGIN GLOBAL FUNCTION DEF: foo
.visible .func (.param.u32 %value_out) foo
{
        .reg.u32 %value;
        .reg.u32 %r22;
        .reg.u32 %r23;
                mov.u32 %r22, 42;
                mov.u32 %r23, %r22;
                mov.u32 %value, %r23;
        st.param.u32    [%value_out], %value;
        ret;
}
...
call (%value_in), bar;
------

Namely: 'bar' only appears in the call.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[burnus at gcc dot gnu.org]

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

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #0)
> .visible .func (.param.u32 %value_out) foo

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#kernel-and-function-directives-alias

has the example:

.visible .func bar(.param .u32 p);
.alias bar, foo;

And the note:
.alias directive requires sm_30 or higher.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[burnus at gcc dot gnu.org]

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

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Created attachment 49239
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=49239&action=edit
Draft patch

PTX ISA Notes
.alias directive introduced in PTX ISA 6.3.

Thus, it does not work everywhere :-(

Patch – I did bump '.version' - but that probably has more side effects than
desirable.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[burnus at gcc dot gnu.org]

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

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #2)
> PTX ISA Notes
> .alias directive introduced in PTX ISA 6.3.

Which is CUDA 10.0, cf.
https://docs.nvidia.com/cuda/archive/10.0/cuda-toolkit-release-notes/index.html#cuda-general-new-features

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #2)
> Created attachment 49239 [details]
> Draft patch
> 
> PTX ISA Notes
> .alias directive introduced in PTX ISA 6.3.
> 
> Thus, it does not work everywhere :-(
> 
> Patch – I did bump '.version' - but that probably has more side effects than
> desirable.

This patch does manage to complete a build, but tests (both standalone and
libgomp) using alias fail with "unresolved symbol".  This is because it calls
write_fn_proto which generates a linker marker for the alias, which makes the
linker think there's such a function, and when it cannot find it, it generates
an "unresolved symbol" error.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 49252
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=49252&action=edit
Other draft patch

I started out independently, and converged to roughly the same code.

One thing I came across during standalone testing is that the alias support in
gcc is for symbols, and as such supports functions and variables, but the ptx
support is only there for functions.  I handle this in the patch with an error,
but it causes quite a few failures:
...
src/gcc/testsuite/gcc.c-torture/execute/alias-2.c:3:12: error: non-function
alias definitions not supported in this configuration^M
compiler exited with status 1
FAIL: gcc.c-torture/execute/alias-2.c   -O0  (test for excess errors)
...
So, that would probably need to be handled by separate effective targets
alias_var/alias_fn.

With standalone testing we also run into various problems related to bumping
the ISA, see also PR96005.

As for libgomp testing, I've tried the test reported in this PR, and got this
ptx link error (silent with trunk, needed to add some printing in plugin):
...
Internal error: reference to deleted section
...

By dumping the ptx code using GOMP_NVPTX_PTXRW and compiling it using ptxas, it
becomes clear that the most complained about problem is deprecated shfl:
...
ptxas gomp-nvptx.11.ptx, line 117; warning : Instruction 'shfl' without '.sync'
is deprecated since PTX ISA version 6.0 and will be discontinued in a future
PTX ISA version
...

The internal error is a driver problem, but it's possible that we push the
driver into a corner by using deprecated insns alongside with modern constructs
like .alias.

So, it seems sensible to do the porting to newer ISA (PR96005) first, before
trying to enable .alias.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
  Attachment #49252|0                           |1
        is obsolete|                            |

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 49253
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=49253&action=edit
Updated other draft patch

(In reply to Tom de Vries from comment #5)
> One thing I came across during standalone testing is that the alias support
> in gcc is for symbols, and as such supports functions and variables, but the
> ptx support is only there for functions.  I handle this in the patch with an
> error, 

Oops, that bit was missing, now added.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Severity|normal                      |enhancement

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

--- Comment #7 from Tom de Vries <vries at gcc dot gnu.org> ---
In PR97106 comment 1, it's suggested:
...
14:11 < amonakov> Tobias__: I think the proper way to solve this is define
hooks for the backend to print something for aliases, and then have nvptx-ld.c
resolve them
...
which wouldn't rely on ptx .alias.

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=97102
Bug 97102 depends on bug 96005, which changed state.

Bug 96005 Summary: Add possibility to use newer ptx isa
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96005

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |RESOLVED
         Resolution|---                         |FIXED

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |12.0
         Resolution|---                         |FIXED
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #8 from Tom de Vries <vries at gcc dot gnu.org> ---
The test-case from comment 0 now works in combination with:
/* { dg-additional-options "-foffload=-malias -foffload=-mptx=6.3" } */

[Bug target/97102] [nvptx] PTX JIT compilation failed when using aliases

[burnus at gcc dot gnu.org]

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

--- Comment #9 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Regarding the commit:

commit r12-7766-gf8b15e177155960017ac0c5daef8780d1127f91c
Author: Tom de Vries <tdevries@suse.de>
Date:   Fri Mar 11 13:41:01 2022 +0100

    [nvptx] Use .alias directive for mptx >= 6.3