https://gcc.gnu.org/g:365437aa677e791b3a1b4a7c6c24a14d551745bd
commit 365437aa677e791b3a1b4a7c6c24a14d551745bd
Author: Prathamesh Kulkarni <prathame...@nvidia.com>
Date: Tue Sep 24 08:18:48 2024 +0530
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 <prathame...@nvidia.com>
Co-authored-by: Thomas Schwinge <tschwi...@baylibre.com>
(cherry picked from commit f5ee372b012594830f6d5f7f4b7e01bae810b1da)
Diff:
---
gcc/ChangeLog.omp | 8 ++++++++
gcc/config/nvptx/nvptx.cc | 24 ++++++++++++++++++++---
gcc/testsuite/ChangeLog.omp | 7 +++++++
gcc/testsuite/gcc.target/nvptx/alias-to-alias-1.c | 6 ++++--
4 files changed, 40 insertions(+), 5 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index dadf7e7aadf8..6077e6bd0fd1 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,5 +1,13 @@
2025-02-18 Thomas Schwinge <tschwi...@baylibre.com>
+ Backported from trunk:
+ 2024-09-24 Prathamesh Kulkarni <prathame...@nvidia.com>
+ Thomas Schwinge <tschwi...@baylibre.com>
+
+ PR target/104957
+ * config/nvptx/nvptx.cc (nvptx_asm_output_def_from_decls): Use
+ cgraph_node::get(name)->ultimate_alias_target instead of value.
+
Backported from trunk:
2024-09-05 Thomas Schwinge <tschwi...@baylibre.com>
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index cfc0a56eae0b..3a55f6d5d64a 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -8006,7 +8006,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)
{
@@ -8041,7 +8042,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;
@@ -8050,11 +8052,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 6a53a73ba1f5..e652800ea46b 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,5 +1,12 @@
2025-02-18 Thomas Schwinge <tho...@codesourcery.com>
+ Backported from trunk:
+ 2024-09-24 Prathamesh Kulkarni <prathame...@nvidia.com>
+ Thomas Schwinge <tschwi...@baylibre.com>
+
+ PR target/104957
+ * gcc.target/nvptx/alias-to-alias-1.c: Adjust.
+
Backported from trunk:
2024-09-05 Thomas Schwinge <tschwi...@baylibre.com>
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 7bce7a358c79..08de9e6d69da 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 } }
