diff mbox

[gomp4] Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs (was: signed nums are better for dimensions)

Message ID 871tf8kdrz.fsf@kepler.schwinge.homeip.net
State New
Headers show

Commit Message

Thomas Schwinge Aug. 12, 2015, 10:21 a.m. UTC 
Hi!

On Tue, 11 Aug 2015 13:38:34 -0400, Nathan Sidwell <nathan@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(-)



Grüße,
 Thomas

Comments

Nathan Sidwell Aug. 12, 2015, 12:46 p.m. UTC | #1 
On 08/12/15 06:21, Thomas Schwinge wrote:
> Hi!
  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

The right solution is probably for the default validated_dims to set the size to 
unity, just like the host.  Thereby forcing any backend that wants to provide a 
larger size to override that hook.

I.e. remove the #ifndef ACCEL_COMPILER from default_validate_dims in omp-low

nathan
Nathan Sidwell Aug. 12, 2015, 12:49 p.m. UTC | #2 
On 08/12/15 08:46, Nathan Sidwell wrote:
> On 08/12/15 06:21, Thomas Schwinge wrote:
>> Hi!
>   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
>
> The right solution is probably for the default validated_dims to set the size to
> unity, just like the host.  Thereby forcing any backend that wants to provide a
> larger size to override that hook.
>
> I.e. remove the #ifndef ACCEL_COMPILER from default_validate_dims in omp-low

for avoidance of doubt, I'll add that to the patch I have in progress.

nathan
diff mbox

Patch

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  <thomas@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  <nathan@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
 }