[Bug target/105019] New: [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[Bug target/105019] New: [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

            Bug ID: 105019
           Summary: [nvptx] malias in libgomp results in "Internal error:
                    reference to deleted section"
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

As mentioned in the commit message for malias:
...
When enabling malias by default, libgomp detects alias support and
consequently libgomp.a will contains a few uses of .alias.  This however
results in aforementioned "Internal error: reference to deleted section" in
many test-cases.  Either there's some error with how .alias is used, or
there's a driver bug.  While this issue is not resolved, we keep malias
off-by-default.
...

This needs to be investigated, and if it's a driver bug, reported to nvidia, or
otherwise fixed or worked around.

Note: the same error showed up in a test-case where the call to an alias was
inlined, and consequently the alias referenced a defined but unused function. 
To observe this, disable this bit:
...
  if (!cgraph_node::get (name)->referred_to_p ())
    /* Prevent "Internal error: reference to deleted section".  */
    return;
...
in nvptx_asm_output_def_from_decls and run nvptx.exp=alias-2.c:
...
PASS: gcc.target/nvptx/alias-2.c (test for excess errors)
spawn nvptx-none-run ./alias-2.exe^M
fatal   : Internal error: reference to deleted section^M
nvptx-run: cuLinkComplete failed: unknown error (CUDA_ERROR_UNKNOWN, 999)^M
FAIL: gcc.target/nvptx/alias-2.c execution test
...

So it's possible that the error somehow related to this scenario.

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
To trigger:
...
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 87efc23bd96..8bf9ea90a77 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -245,6 +245,9 @@ default_ptx_version_option (void)
      warp convergence.  */
   res = MAX (res, PTX_VERSION_6_0);

+  /* Pick at least 6.3, to enable using malias.  */
+  res = MAX (res, PTX_VERSION_6_3);
+
   /* Verify that we pick a version that supports the sm.  */
   gcc_assert (first <= res);
   return res;
diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index 11288d1a8ee..a4aece80682 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -87,7 +87,7 @@ mptx-comment
 Target Var(nvptx_comment) Init(1) Undocumented

 malias-
-Target Var(nvptx_alias) Init(0) Undocumented
+Target Var(nvptx_alias) Init(1) Undocumented

 mexperimental
 Target Var(nvptx_experimental) Init(0) Undocumented
...
rebuild gcc, run libgomp tests, and:
...
$ grep -c "Internal error: reference to deleted section" libgomp.log 
637
...

For instance:
...
Execution timeout is: 300
spawn [open ...]^M

libgomp: Link error log fatal   : Internal error: reference to deleted section


libgomp: cuLinkComplete error: unknown error

libgomp: Cannot map target functions or variables (expected 2, have 4294967295)
FAIL: libgomp.c/../libgomp.c-c++-common/declare_target-1.c execution test
...

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Aliases in libgomp.a:
...
$ grep "\.alias"
build-gcc-offload-nvptx-none/nvptx-none/mgomp/libgomp/.libs/libgomp.a
.alias gomp_ialias_GOMP_loop_runtime_next,GOMP_loop_runtime_next;
.alias gomp_ialias_GOMP_loop_ull_runtime_next,GOMP_loop_ull_runtime_next;
.alias gomp_ialias_GOMP_parallel_end,GOMP_parallel_end;
.alias gomp_ialias_GOMP_taskgroup_start,GOMP_taskgroup_start;
.alias gomp_ialias_GOMP_taskgroup_end,GOMP_taskgroup_end;
.alias
gomp_ialias_GOMP_taskgroup_reduction_register,GOMP_taskgroup_reduction_register;
.alias gomp_ialias_omp_capture_affinity,omp_capture_affinity;
.alias gomp_ialias_omp_aligned_alloc,omp_aligned_alloc;
.alias gomp_ialias_omp_free,omp_free;
.alias gomp_ialias_omp_aligned_calloc,omp_aligned_calloc;
...

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
Aliases in failing .exe:
...
$ strings declare_target-1.exe | grep "\.alias"
.alias gomp_ialias_GOMP_taskgroup_start,GOMP_taskgroup_start;
.alias gomp_ialias_GOMP_taskgroup_end,GOMP_taskgroup_end;
.alias
gomp_ialias_GOMP_taskgroup_reduction_register,GOMP_taskgroup_reduction_register;
...

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
OK, I think this is the pattern:
...
$ cat gcc/testsuite/gcc.target/nvptx/alias-5.c
/* { dg-do link } */
/* { dg-do run { target runtime_ptx_isa_version_6_3 } } */
/* { dg-options "-save-temps -malias -mptx=6.3" } */

int v;

void __attribute__((noinline))
__f ()
{
  v = 1;
}

void f () __attribute__ ((alias ("__f")));

void
g (void)
{
  f ();
}

int
main (void)
{
  if (v != 0)
    __builtin_abort ();

  return 0;
}
...

There's a function __f:
...
// BEGIN GLOBAL FUNCTION DEF: __f
.visible .func __f
{
.reg .u32 %r22;
mov.u32 %r22,1;
st.global.u32 [v],%r22;
ret;
}
...
with alias f:
...
// BEGIN GLOBAL FUNCTION DECL: __f
.visible .func __f;
.visible .func f;
.alias f,__f;
...
called from g:
...
// BEGIN GLOBAL FUNCTION DEF: g
.visible .func g
{
{
call f;
}
ret;
}
...

However, g is unused.

So we have:
...
PASS: gcc.target/nvptx/alias-5.c (test for excess errors)
spawn nvptx-none-run ./alias-5.exe^M
fatal   : Internal error: reference to deleted section^M
nvptx-run: cuLinkComplete failed: unknown error (CUDA_ERROR_UNKNOWN, 999)^M
FAIL: gcc.target/nvptx/alias-5.c execution test
...

Calling g from main fixes the internal error.

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
Creating a CUDA example is hampered by the fact that there's no symbol alias
support, AFAICT.

I'd like to write something like:
...
__device__ void __foo ()
{
  printf ("__foo\n");
}
__device__ void foo () __attribute__((alias ("__foo")));
__device__ void bar ()
{
  foo ();
}
__global__ void hello_world ()
{
  bar ();
}
...
(and then comment out bar to reproduce the problem), but all attempts to get
this compiled and executed end up in various errors.

Of course we can resort to hand-edited ptx, or adding .alias directives in asm
statements, but that's likely to produce an 'unsupported' response by nvidia
when filed as bug report.

[Bug target/105019] [nvptx] malias in libgomp results in "Internal error: reference to deleted section"

[vries at gcc dot gnu.org]

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

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #4)
> OK, I think this is the pattern:
> ...
> $ cat gcc/testsuite/gcc.target/nvptx/alias-5.c

FTR, same thing if I use static functions:
...
static void __attribute__((noinline))
__f ()
{
  v = 1;
}

static void f () __attribute__ ((alias ("__f")));

static void
g (void)
{
  f ();
}
...