nvptx: Partial support for aliases to aliases.

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;
}

gcc emits following ptx for baz:
.visible .func (.param.u32 %value_out) bar;
.alias bar,foo;
.visible .func (.param.u32 %value_out) baz;
.alias baz,bar;

which is incorrect since PTX requires aliasee to be a defined function.
The patch instead uses cgraph_node::get(name)->ultimate_alias_target,
which generates the following PTX:

.visible .func (.param.u32 %value_out) baz;
.alias baz,foo;

gcc/ChangeLog:
	PR target/104957
	* config/nvptx/nvptx.cc (nvptx_asm_output_def_from_decls): Use
	cgraph_node::get(name)->ultimate_alias_target instead of value.

gcc/testsuite/ChangeLog:
	PR target/104957
	* gcc.target/nvptx/alias-to-alias-1.c: Adjust.

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>

2 files changed, 25 insertions, 5 deletions

diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 4a7c64f..96a1134 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -7582,7 +7582,8 @@ nvptx_mem_local_p (rtx mem)
   while (0)
 
 void
-nvptx_asm_output_def_from_decls (FILE *stream, tree name, tree value)
+nvptx_asm_output_def_from_decls (FILE *stream, tree name,
+				 tree value ATTRIBUTE_UNUSED)
 {
   if (nvptx_alias == 0 || !TARGET_PTX_6_3)
     {
@@ -7617,7 +7618,8 @@ nvptx_asm_output_def_from_decls (FILE *stream, tree name, tree value)
       return;
     }
 
-  if (!cgraph_node::get (name)->referred_to_p ())
+  cgraph_node *cnode = cgraph_node::get (name);
+  if (!cnode->referred_to_p ())
     /* Prevent "Internal error: reference to deleted section".  */
     return;
 
@@ -7626,11 +7628,27 @@ nvptx_asm_output_def_from_decls (FILE *stream, tree name, tree value)
   fputs (s.str ().c_str (), stream);
 
   tree id = DECL_ASSEMBLER_NAME (name);
+
+  /* Walk alias chain to get reference callgraph node.
+     The rationale of using ultimate_alias_target here is that
+     PTX's .alias directive only supports 1-level aliasing where
+     aliasee is function defined in same module.
+
+     So for the following case:
+     int foo() { return 42; }
+     int bar () __attribute__((alias ("foo")));
+     int baz () __attribute__((alias ("bar")));
+
+     should resolve baz to foo:
+     .visible .func (.param.u32 %value_out) baz;
+     .alias baz,foo;  */
+  symtab_node *alias_target_node = cnode->ultimate_alias_target ();
+  tree alias_target_id = DECL_ASSEMBLER_NAME (alias_target_node->decl);
   std::stringstream s_def;
   write_fn_marker (s_def, true, TREE_PUBLIC (name), IDENTIFIER_POINTER (id));
   fputs (s_def.str ().c_str (), stream);
   NVPTX_ASM_OUTPUT_DEF (stream, IDENTIFIER_POINTER (id),
-			IDENTIFIER_POINTER (value));
+			IDENTIFIER_POINTER (alias_target_id));
 }
 
 #undef NVPTX_ASM_OUTPUT_DEF
diff --git a/gcc/testsuite/gcc.target/nvptx/alias-to-alias-1.c b/gcc/testsuite/gcc.target/nvptx/alias-to-alias-1.c
index 7bce7a3..08de9e6 100644
--- a/gcc/testsuite/gcc.target/nvptx/alias-to-alias-1.c
+++ b/gcc/testsuite/gcc.target/nvptx/alias-to-alias-1.c
@@ -1,6 +1,8 @@
 /* Alias to alias; 'libgomp.c-c++-common/pr96390.c'.  */
 
-/* { dg-do compile } */
+/* { dg-do link } */
+/* { dg-do run { target nvptx_runtime_alias_ptx } } */
+/* { dg-options -save-temps } */
 /* { dg-add-options nvptx_alias_ptx } */
 
 int v;
@@ -32,7 +34,7 @@ main (void)
 /* { dg-final { scan-assembler-times {(?n)^// BEGIN GLOBAL FUNCTION DECL: baz$} 1 } }
    { dg-final { scan-assembler-times {(?n)^\.visible \.func baz;$} 1 } }
    { dg-final { scan-assembler-times {(?n)^// BEGIN GLOBAL FUNCTION DEF: baz$} 1 } }
-   { dg-final { scan-assembler-times {(?n)^\.alias baz,bar;$} 1 } } */
+   { dg-final { scan-assembler-times {(?n)^\.alias baz,foo;$} 1 } } */
 
 /* { dg-final { scan-assembler-times {(?n)\tcall foo;$} 0 } }
    { dg-final { scan-assembler-times {(?n)\tcall bar;$} 0 } }