[Bug target/96991] New: [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[Bug target/96991] New: [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[vries at gcc dot gnu.org]

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

            Bug ID: 96991
           Summary: [nvptx] internal compiler error: in write_fn_proto, at
                    config/nvptx/nvptx.c:913
           Product: gcc
           Version: 11.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: ---

When running this libgomp testsuite test-case on x86_64 with nvptx accelerator:
...
$ cat src/libgomp/testsuite/libgomp.c/atomic.c
/* { dg-do run } */

#include <stdbool.h>

__uint128_t v;
#pragma omp declare target (v)

int
main ()
{
  #pragma omp target
  {
    __atomic_add_fetch (&v, 1, __ATOMIC_RELAXED);
    __atomic_fetch_add (&v, 1, __ATOMIC_RELAXED);
    __uint128_t exp = 2;
    __atomic_compare_exchange_n (&v, &exp, 7, false, __ATOMIC_RELEASE,
__ATOMIC_ACQUIRE);
  }
} 
...
we run into an assert:
...
/home/vries/oacc/trunk/source-gcc/libgomp/testsuite/libgomp.c/atomic.c: In
function ‘main._omp_fn.0’:
/home/vries/oacc/trunk/source-gcc/libgomp/testsuite/libgomp.c/atomic.c:11:9:
internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913
0x15ffdef write_fn_proto
        /home/vries/oacc/trunk/source-gcc/gcc/config/nvptx/nvptx.c:913
0x160039c nvptx_record_fndecl
        /home/vries/oacc/trunk/source-gcc/gcc/config/nvptx/nvptx.c:989
0x160676d nvptx_output_call_insn(rtx_insn*, rtx_def*, rtx_def*)
        /home/vries/oacc/trunk/source-gcc/gcc/config/nvptx/nvptx.c:2417
0x1a91699 output_150
        /home/vries/oacc/trunk/source-gcc/gcc/config/nvptx/nvptx.md:887
0xbae099 get_insn_template(int, rtx_insn*)
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:2070
0xbb0af9 final_scan_insn_1
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:3039
0xbb125f final_scan_insn(rtx_insn*, _IO_FILE*, int, int, int*)
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:3152
0xbade27 final_1
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:2020
0xbb636e rest_of_handle_final
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:4658
0xbb683e execute
        /home/vries/oacc/trunk/source-gcc/gcc/final.c:4736
Please submit a full bug report,
with preprocessed source if appropriate.
Please include the complete backtrace with any bug report.
See <https://gcc.gnu.org/bugs/> for instructions.
...

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
We run into assert:
...
913             gcc_assert (type == boolean_type_node);
...
because:
...
(gdb) call debug_generic_expr (type)
_Bool
(gdb) call debug_generic_expr (boolean_type_node)
No symbol "boolean_type_node" in current context.
(gdb) call debug_generic_expr (global_trees[TI_BOOLEAN_TYPE])
bool
...

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[vries at gcc dot gnu.org]

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

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Patch:
...
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 39d0275493a..6f393dfea01 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -910,7 +910,7 @@ write_fn_proto (std::stringstream &s, bool is_defn,
       if (not_atomic_weak_arg)
        argno = write_arg_type (s, -1, argno, type, prototyped);
       else
-       gcc_assert (type == boolean_type_node);
+       gcc_assert (TREE_CODE (type) == BOOLEAN_TYPE);
     }

   if (stdarg_p (fntype))
...

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jakub at gcc dot gnu.org

--- Comment #3 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Note, the testcase would need to require int128 effective target or even
sync_int_128_runtime, ensure linking with -latomic on offloading targets that
need it (unless we handle it automatically), but also as I said, it contains
just 3 randomly chosen atomics.
Would be nice to have more complete test coverage of the atomics, like having
one test per basic atomic size (char, short, int, long long, int128) and in
each test at least most of the __atomic* and __sync_* builtins that are out
there.
See e.g. gcc.dg/atomic*.c and gcc.dg/sync*.c.  Perhaps use macros to avoid too
many repetitions.

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Jakub Jelinek from comment #3)
> Note, the testcase would need to require int128 effective target or even
> sync_int_128_runtime, ensure linking with -latomic on offloading targets
> that need it (unless we handle it automatically), but also as I said, it
> contains just 3 randomly chosen atomics.
> Would be nice to have more complete test coverage of the atomics, like
> having one test per basic atomic size (char, short, int, long long, int128)
> and in each test at least most of the __atomic* and __sync_* builtins that
> are out there.
> See e.g. gcc.dg/atomic*.c and gcc.dg/sync*.c.  Perhaps use macros to avoid
> too many repetitions.

For now, I'm going to commit without test-case.

I've tried for a while to make something that does not depend on libatomic, but
didn't manage.

And I didn't manage to xfail the resulting link failure due to no libatomic
support in a way that doesn't also xfail the ICE.

So, to be revisited once libatomic support for nvptx has been committed.

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[cvs-commit at gcc dot gnu.org]

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

--- Comment #5 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:505590b796df18ec3fcdcd6b8060f6f1410660b2

commit r11-3073-g505590b796df18ec3fcdcd6b8060f6f1410660b2
Author: Tom de Vries <tdevries@suse.de>
Date:   Wed Sep 9 09:51:43 2020 +0200

    [nvptx] Fix boolean type test in write_fn_proto

    When running this libgomp testcase for nvptx accelerator:
    ...
    /* { dg-do run } */
    __uint128_t v;
    int main () {
      #pragma omp target
      {
        __uint128_t exp = 2;
        __atomic_compare_exchange_n (&v, &exp, 7, false, __ATOMIC_RELEASE,
                                     __ATOMIC_ACQUIRE);
      }
    }
    ...
    we run into this assert in write_fn_proto:
    ...
    913             gcc_assert (type == boolean_type_node);
    ...

    This happens when doing some special-handling code for
    __atomic_compare_exchange_1/2/4/8/16.  The function decls have a parameter
    called weak of type bool, which is skipped when writing the decl because
    the corresponding libatomic functions do not have that parameter.  The
assert
    is there to verify that we skip the correct parameter.

    However, we assert because we have different type of bools:
    ...
    (gdb) call debug_generic_expr (type)
    _Bool
    (gdb) call debug_generic_expr (global_trees[TI_BOOLEAN_TYPE])
    bool
    ...

    Fix this by checking for TREE_CODE (type) == BOOLEAN_TYPE instead.

    Tested libgomp on x86_64-linux with nvptx accelerator.

    Likewise, tested that the test-case above does not ICE anymore.

    gcc/ChangeLog:

            PR target/96991
            * config/nvptx/nvptx.c (write_fn_proto): Fix boolean type check.

[Bug target/96991] [nvptx] internal compiler error: in write_fn_proto, at config/nvptx/nvptx.c:913

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
             Status|UNCONFIRMED                 |RESOLVED
   Target Milestone|---                         |11.0

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Patch committed, test-case postponed for when there's libatomic support for
nvptx.

Marking resolved-fixed.