Hi! On Tue, 11 Aug 2015 13:38:34 -0400, Nathan Sidwell <nat...@acm.org> wrote: > 2) We really should not be getting to the expanders if there's nothing to > expand > to. That's simply covering up lack of earlier handling. That earlier > removal > gives optimizers more leeway.
> --- internal-fn.c (revision 226770)
> +++ internal-fn.c (working copy)
> static void
> -expand_GOACC_FORK (gcall *stmt ATTRIBUTE_UNUSED)
> +expand_GOACC_FORK (gcall *ARG_UNUSED (stmt))
> {
> #ifdef HAVE_oacc_fork
> [...]
> +#else
> + gcc_unreachable ();
> #endif
> }
>
> static void
> -expand_GOACC_JOIN (gcall *stmt ATTRIBUTE_UNUSED)
> +expand_GOACC_JOIN (gcall *ARG_UNUSED (stmt))
> {
> #ifdef HAVE_oacc_join
> [...]
> +#else
> + gcc_unreachable ();
> #endif
> }
> static void
> -expand_GOACC_DIM_SIZE (gcall *stmt)
> +expand_GOACC_DIM_SIZE (gcall *ARG_UNUSED (stmt))
> {
> +#ifdef HAVE_oacc_dim_size
> [...]
> #else
> - emit_move_insn (target, const1_rtx);
> + gcc_unreachable ();
> #endif
> }
>
> static void
> -expand_GOACC_DIM_POS (gcall *stmt)
> +expand_GOACC_DIM_POS (gcall *ARG_UNUSED (stmt))
> {
> +#ifdef HAVE_oacc_dim_pos
> [...]
> #else
> - emit_move_insn (target, const0_rtx);
> + gcc_unreachable ();
> #endif
> }
It's not an issue for expand_GOACC_FORK and expand_GOACC_JOIN, but
expand_GOACC_DIM_SIZE and expand_GOACC_DIM_POS then blow up when the
intelmic offloading compiler works through the OpenACC offloading code,
for example:
[...]/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/firstprivate-1.c:
In function 'main._omp_fn.0':
[...]/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/firstprivate-1.c:18:11:
internal compiler error: in expand_GOACC_DIM_POS, at internal-fn.c:2023
#pragma acc parallel num_gangs (n) firstprivate (a)
^
0x90bce7 expand_GOACC_DIM_POS
[...]/source-gcc/gcc/internal-fn.c:2023
0x62945a expand_call_stmt
[...]/source-gcc/gcc/cfgexpand.c:2279
0x62945a expand_gimple_stmt_1
[...]/source-gcc/gcc/cfgexpand.c:3238
0x62945a expand_gimple_stmt
[...]/source-gcc/gcc/cfgexpand.c:3399
0x62a82d expand_gimple_basic_block
[...]/source-gcc/gcc/cfgexpand.c:5411
0x62fc86 execute
[...]/source-gcc/gcc/cfgexpand.c:6023
Please submit a full bug report,
with preprocessed source if appropriate.
Please include the complete backtrace with any bug report.
See <http://gcc.gnu.org/bugs.html> for instructions.
mkoffload-intelmic: fatal error:
[...]/install/offload-x86_64-intelmicemul-linux-gnu/bin//x86_64-unknown-linux-gnu-accel-x86_64-intelmicemul-linux-gnu-gcc
returned 1 exit status
Admittedly, compiling OpenACC offloading code for Intel MIC doesn't make
a lot of sense (currently), but it is what's being done, and has caused a
lot of regressions in my testing, so I committed the following workaround
to gomp-4_0-branch in r226804. A different approach would have been to
invent some machinery to not compile OpenACC offloading code on the Intel
MIC offloading path, but as the latter eventually is to support OpenACC
offloading, too (see
<http://news.gmane.org/find-root.php?message_id=%3C20141112111207.GW5026%40tucnak.redhat.com%3E>,
for example), this seemed like a step into the wrong direction.
Eventually, this will need to be fixed differently/properly -- have to
implement oacc_dim_size and oacc_dim_pos for Intel MIC offloading, or can
we have a generic fallback solution not specific to the offloading
target?
commit 8a3187f17e13bd45e630ff8a587c1fc7086abece
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Aug 12 09:56:59 2015 +0000
Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs
Followup to r226783.
gcc/
* internal-fn.c: Include "builtins.h".
(expand_GOACC_DIM_SIZE, expand_GOACC_DIM_POS): Instead of
gcc_unreachable, expand_builtin_trap.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@226804
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 6 ++++++
gcc/internal-fn.c | 5 +++--
2 files changed, 9 insertions(+), 2 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index ea254c5..044fdf7 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2015-08-12 Thomas Schwinge <tho...@codesourcery.com>
+
+ * internal-fn.c: Include "builtins.h".
+ (expand_GOACC_DIM_SIZE, expand_GOACC_DIM_POS): Instead of
+ gcc_unreachable, expand_builtin_trap.
+
2015-08-10 Nathan Sidwell <nat...@acm.org>
* tree-ssa-phiopt.c (minmax_replacement): Create new ssa name if
diff --git gcc/internal-fn.c gcc/internal-fn.c
index 8b8c6e1..70bffd4 100644
--- gcc/internal-fn.c
+++ gcc/internal-fn.c
@@ -47,6 +47,7 @@ along with GCC; see the file COPYING3. If not see
#include "stringpool.h"
#include "tree-ssanames.h"
#include "diagnostic-core.h"
+#include "builtins.h"
/* The names of each internal function, indexed by function number. */
const char *const internal_fn_name_array[] = {
@@ -2003,7 +2004,7 @@ expand_GOACC_DIM_SIZE (gcall *ARG_UNUSED (stmt))
VOIDmode, EXPAND_NORMAL);
emit_insn (gen_oacc_dim_size (target, dim));
#else
- gcc_unreachable ();
+ expand_builtin_trap ();
#endif
}
@@ -2021,7 +2022,7 @@ expand_GOACC_DIM_POS (gcall *ARG_UNUSED (stmt))
VOIDmode, EXPAND_NORMAL);
emit_insn (gen_oacc_dim_pos (target, dim));
#else
- gcc_unreachable ();
+ expand_builtin_trap ();
#endif
}
Grüße,
Thomas
signature.asc
Description: PGP signature
