[Bug middle-end/94233] New: [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[Bug middle-end/94233] New: [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[burnus at gcc dot gnu.org]

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

            Bug ID: 94233
           Summary: [OpenMP] ICE in lto1 for AMDGCN offloading: in
                    dr_analyze_innermost, at tree-data-ref.c:925 for
                    gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1
                    .c
           Product: gcc
           Version: 10.0
            Status: UNCONFIRMED
          Keywords: ice-on-valid-code, openmp
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

Compiling gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c with
"-fopenmp -O3" (instead of default -O2) fails as follows.

[I also did see "libgomp: Cannot map target variables (size mismatch)" at run
time for -O2.]


The ICE is in the offloading-target compiler lto1:


/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/install/libexec/gcc/x86_64-none-linux-gnu/10.0.1/accel/amdgcn-amdhsa/lto1
-quiet -dumpbase ccFoh0bA.o -m64 -mgomp -mlocal-symbol-id=/tmp/ccFoh0bA.o
-march=fiji -auxbase-strip /tmp/ccxUIhxG.mkoffload.1.s -g -O3 -version
-fno-openacc -fno-pie -foffload-abi=lp64 -fopenmp -o
/tmp/ccxUIhxG.mkoffload.1.s @/tmp/cciArjCH
GNU GIMPLE (GCC) version 10.0.1 20200319 (experimental) (amdgcn-amdhsa)
        compiled by GNU C version 8.3.1 20190416, GMP version 6.1.2, MPFR
version 4.0.2-p1, MPC version 1.1.0, isl version isl-0.20-GMP
...

src/gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c: In function
'main._omp_fn.0':
src/gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c:38:11: internal
compiler error: in dr_analyze_innermost, at tree-data-ref.c:925
   38 |   #pragma omp target map (to: a, b, c, d) map (from: res)
      |           ^
0x7b8c81 dr_analyze_innermost(innermost_loop_behavior*, tree_node*, loop*,
gimple const*)
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-data-ref.c:924
0x13da92b create_data_ref(edge_def*, loop*, tree_node*, gimple*, bool, bool)
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-data-ref.c:1255
0x13dacb7 find_data_references_in_stmt(loop*, gimple*, vec<data_reference*,
va_heap, vl_ptr>*)
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-data-ref.c:5603
0x13f3093 vect_find_stmt_data_reference(loop*, gimple*, vec<data_reference*,
va_heap, vl_ptr>*)
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-vect-data-refs.c:4064
0xf1f2c6 vect_slp_bb(basic_block_def*)
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-vect-slp.c:3499
0xf2257f execute
       
/net/build5-trusty-cs/scratch/tburnus/fsf.mainline.x86_64-linux-gnu-amdgcn/src/gcc-mainline/gcc/tree-vectorizer.c:1321

[Bug middle-end/94233] [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[rguenth at gcc dot gnu.org]

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

--- Comment #1 from Richard Biener <rguenth at gcc dot gnu.org> ---
  get_object_alignment_1 (base, &bit_base_alignment, &bit_base_misalignment);

  /* There are no bitfield references remaining in BASE, so the values
     we got back must be whole bytes.  */
  gcc_assert (bit_base_alignment % BITS_PER_UNIT == 0
              && bit_base_misalignment % BITS_PER_UNIT == 0);

where base is the result of get_inner_reference.  This possibly hints at
some odd alignment values on some decl(?).  Needs debugging of the
actual base seen here and how it transfers from compile to offload lto1.

[Bug middle-end/94233] [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |ASSIGNED
     Ever confirmed|0                           |1
   Last reconfirmed|                            |2020-03-20

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Richard Biener from comment #1)
>   get_object_alignment_1 (base, &bit_base_alignment, &bit_base_misalignment);
>   gcc_assert (bit_base_alignment % BITS_PER_UNIT == 0
>               && bit_base_misalignment % BITS_PER_UNIT == 0);
> where base is the result of get_inner_reference.  This possibly hints at
> some odd alignment values on some decl(?).

Indeed, bit_base_alignment is here == 1; the alignment is:
  ((unsigned)1) << ((NODE)->decl_common.align - 1
hence,  …->decl_common.align == 1

The align:1 is produced by lto.c's offload_handle_link_vars(), which calls:
  tree link_ptr_var = make_node (VAR_DECL);
which contains:
  enum tree_code_class type = TREE_CODE_CLASS (code);
...
  switch (type)
...
    case tcc_declaration:
...
            SET_DECL_ALIGN (t, 1);


Untested patch:

diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 39bb5f45c95..812dcc39b44 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -568,2 +568,3 @@ offload_handle_link_vars (void)
        SET_DECL_MODE (link_ptr_var, TYPE_MODE (type));
+       SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
        DECL_SIZE (link_ptr_var) = TYPE_SIZE (type);

[Bug middle-end/94233] [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[cvs-commit at gcc dot gnu.org]

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

--- Comment #3 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tobias Burnus <burnus@gcc.gnu.org>:

https://gcc.gnu.org/g:b809f0b6580969c4f047f4dae072c090718efd76

commit r10-7326-gb809f0b6580969c4f047f4dae072c090718efd76
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Mon Mar 23 08:31:23 2020 +0100

    Set proper DECL_ALIGN in offload_handle_link_vars (PR94233)

            gcc/lto/
            PR middle-end/94233
            * lto.c (offload_handle_link_vars): Cleanup; call
            build_decl to ensure alignment is set.

[Bug middle-end/94233] [OpenMP] ICE in lto1 for AMDGCN offloading: in dr_analyze_innermost, at tree-data-ref.c:925 for gcc-mainline/libgomp/testsuite/libgomp.c/target-link-1.c

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
             Status|ASSIGNED                    |RESOLVED

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
FIXED for GCC 10 trunk.

For the run-time issue, see PR94251.