> -----Original Message-----
> From: Prathamesh Kulkarni <prathame...@nvidia.com>
> Sent: Monday, August 26, 2024 4:21 PM
> To: Thomas Schwinge <tschwi...@baylibre.com>; gcc-patches@gcc.gnu.org
> Subject: [nvptx] Fix code-gen for alias attribute
>
> External email: Use caution opening links or attachments
>
>
> Hi,
> For the following test (adapted from pr96390.c):
>
> __attribute__((noipa)) int foo () { return 42; } int bar ()
> __attribute__((alias ("foo"))); int baz () __attribute__((alias
> ("bar")));
>
> int main ()
> {
> int n;
> #pragma omp target map(from:n)
> n = baz ();
> return n;
> }
>
> Compiling with -fopenmp -foffload=nvptx-none -foffload=-malias -
> foffload=-mptx=6.3 results in:
>
> ptxas fatal : Internal error: alias to unknown symbol
> nvptx-as: ptxas returned 255 exit status nvptx mkoffload: fatal error:
> ../../install/bin/aarch64-unknown-linux-gnu-accel-nvptx-none-gcc
> returned 1 exit status compilation terminated.
> lto-wrapper: fatal error: /home/prathameshk/gnu-toolchain/gcc/grcogcc-
> 38/install/libexec/gcc/aarch64-unknown-linux-gnu/15.0.0//accel/nvptx-
> none/mkoffload returned 1 exit status compilation terminated.
>
> This happens because ptx code-gen shows:
>
> // BEGIN GLOBAL FUNCTION DEF: foo
> .visible .func (.param.u32 %value_out) foo {
> .reg.u32 %value;
> mov.u32 %value, 42;
> st.param.u32 [%value_out], %value;
> ret;
> }
> .visible .func (.param.u32 %value_out) bar; .alias bar,foo; .visible
> .func (.param.u32 %value_out) baz; .alias baz,bar;
>
> .alias baz, bar is invalid since PTX requires aliasee to be a defined
> function:
> https://sw-docs-dgx-station.nvidia.com/cuda-latest/parallel-thread-
> execution/latest-internal/#kernel-and-function-directives-alias
>
> The patch uses cgraph_node::get(name)->ultimate_alias_target ()
> instead of the provided value in nvptx_asm_output_def_from_decls.
> For the above case, it now generates the following ptx:
>
> .alias baz,foo;
> instead of:
> .alias baz,bar;
>
> which fixes the issue.
>
> Does the patch look in the right direction ?
>
> Signed-off-by: Prathamesh Kulkarni <prathame...@nvidia.com>
Hi,
ping: https://gcc.gnu.org/pipermail/gcc-patches/2024-August/661457.html
Thanks,
Prathamesh
>
> Thanks,
> Prathamesh
