> -----Original Message-----
> From: Thomas Schwinge <tschwi...@baylibre.com>
> Sent: Wednesday, September 4, 2024 3:15 PM
> To: Prathamesh Kulkarni <prathame...@nvidia.com>; Jan Hubicka
> <hubi...@ucw.cz>; gcc-patches@gcc.gnu.org
> Subject: Re: [nvptx] Fix code-gen for alias attribute
>
> External email: Use caution opening links or attachments
>
>
> Hi!
>
> Honza (or others, of course), there's a question about
> 'ultimate_alias_target'.
>
> On 2024-08-26T10:50:36+0000, Prathamesh Kulkarni
> <prathame...@nvidia.com> wrote:
> > For the following test (adapted from pr96390.c):
> >
> > __attribute__((noipa)) int foo () { return 42; } int bar ()
> > __attribute__((alias ("foo"))); int baz () __attribute__((alias
> > ("bar")));
>
> > Compiling [for nvptx] results in:
> >
> > ptxas fatal : Internal error: alias to unknown symbol
> > nvptx-as: ptxas returned 255 exit status
>
> Prathamesh: thanks for looking into this, and ACK: one of the many
> limitations of PTX '.alias'. :-|
>
> > This happens because ptx code-gen shows:
> >
> > // BEGIN GLOBAL FUNCTION DEF: foo
> > .visible .func (.param.u32 %value_out) foo {
> > [...]
> > }
> > .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-exe
> > cution/latest-internal/#kernel-and-function-directives-alias
>
> (Us ordinary mortals need to look at
> <https://docs.nvidia.com/cuda/parallel-thread-execution/#kernel-and-
> function-directives-alias>;
> please update the Git commit log.)
>
> > The patch uses cgraph_node::get(name)->ultimate_alias_target ()
> instead of the provided value in nvptx_asm_output_def_from_decls.
>
> I confirm that resolving to 'ultimate_alias_target' does work for this
> case:
>
> > For the above case, it now generates the following ptx:
> >
> > .alias baz,foo;
> > instead of:
> > .alias baz,bar;
> >
> > which fixes the issue.
>
> ..., but I'm not sure if that's conceptually correct; I'm not familiar
> with 'ultimate_alias_target' semantics. (Honza?)
>
> Also, I wonder whether 'gcc/varasm.cc:do_assemble_alias' is prepared for
> 'ASM_OUTPUT_DEF_FROM_DECLS' to disregard the specified 'target'/'value'
> and instead do its own thing (here, the proposed resolving to
> 'ultimate_alias_target')? (No other GCC back end appears to be doing
> such a thing; from a quick look, all appear to faithfully use the
> specified 'target'/'value'.)
>
> Now, consider the case that the source code is changed as follows:
>
> __attribute__((noipa)) int foo () { return 42; }
> -int bar () __attribute__((alias ("foo")));
> +int bar () __attribute__((weak, alias ("foo")));
> int baz () __attribute__((alias ("bar")));
>
> With 'ultimate_alias_target', I've checked, you'd then still emit
> '.alias baz,foo;', losing the ability to override the weak alias with a
> strong 'bar' definition in another compilation unit?
>
> Now, that said: GCC/nvptx for such code currently diagnoses
> "error: weak alias definitions not supported [...]" ;-| -- so we may be
> safe, after all? ..., or is there any other way that the resolving to
> 'ultimate_alias_target' might cause issues? If not, then at least your
> proposed patch shouldn't be causing any harm (doesn't affect '--
> target=nvptx-none' test results at all...), and does address one user-
> visible issue ('libgomp.c-c++-common/pr96390.c'), and thus makes sense
> to install.
>
> > [nvptx] Fix code-gen for alias attribute.
>
> I'd rather suggest something like:
> "[nvptx] (Some) support for aliases to aliases" (or similar).
>
> Also, please add "PR target/104957" to the Git commit log, as your
> change directly alters this one aspect of PR104957 "[nvptx] Use .alias
> directive (available starting ptx isa version 6.3)"'s commit r12-7766-
> gf8b15e177155960017ac0c5daef8780d1127f91c
> "[nvptx] Use .alias directive for mptx >= 6.3":
>
> | Aliases to aliases are not supported (see libgomp.c-c++-
> common/pr96390.c).
> | This is currently not prohibited by the compiler, but with the driver
> | link we run into: "Internal error: alias to unknown symbol" .
>
> ... which we then have (some) support for with the proposed code
> changes:
>
> > --- a/gcc/config/nvptx/nvptx.cc
> > +++ b/gcc/config/nvptx/nvptx.cc
> > @@ -7583,7 +7583,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)
> > {
> > @@ -7618,7 +7619,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;
> >
> > @@ -7627,8 +7629,10 @@ nvptx_asm_output_def_from_decls (FILE *stream,
> tree name, tree value)
> > fputs (s.str ().c_str (), stream);
> >
> > tree id = DECL_ASSEMBLER_NAME (name);
> > + symtab_node *alias_target_node = cnode->ultimate_alias_target ();
> > + tree alias_target_id = DECL_ASSEMBLER_NAME
> > + (alias_target_node->decl);
> > NVPTX_ASM_OUTPUT_DEF (stream, IDENTIFIER_POINTER (id),
> > - IDENTIFIER_POINTER (value));
> > + IDENTIFIER_POINTER (alias_target_id));
> > }
> >
> > #undef NVPTX_ASM_OUTPUT_DEF
>
> Please put some rationale comment before the 'ultimate_alias_target'.
>
>
> All that said, I'm currently working on <https://gcc.gnu.org/PR105018>
> "[nvptx] Need better alias support", via
> <https://github.com/SourceryTools/nvptx-tools/issues/32>
> "[LD] Handle alias in nvptx-ld as nvptx's .alias does not handle it
> fully".
Hi Thomas,
Thanks for the review and sorry for late reply.
The attached patch addresses the above suggestions.
Does it look OK ?
(Also, could you please test it at your end as well?)
Signed-off-by:
Thanks,
Prathamesh
>
>
> Grüße
> Thomas
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>
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 4a7c64f05eb..96a1134220e 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 7bce7a358c7..08de9e6d69d 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 } }
