Fix decimal floating-point LTO streaming for offloading compilation

Message ID 20191128142402.2949931d@squid.athome
State New
Headers show
Series
  • Fix decimal floating-point LTO streaming for offloading compilation
Related show

Commit Message

Julian Brown Nov. 28, 2019, 2:24 p.m.
As mentioned in PR91985, offloading compilation is broken at present
because of an issue with LTO streaming. With thanks to Joseph for
hints, here's a solution.

Unlike e.g. the _FloatN types, when decimal floating-point types are
enabled, common tree nodes are created for each float type size (e.g.
dfloat32_type_node) and also a pointer to each type is created
(e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node emits
these like:

  <float:32>     (dfloat32_type_node)
  <float:64>     (dfloat64_type_node)
  <float:128>    (dfloat128_type_node)
  <float:32> *   (dfloat32_ptr_type_node)
  <float:32>
  <float:64> *   (dfloat64_ptr_type_node)
  <float:64>
  <float:128> *  (dfloat128_ptr_type_node)
  <float:128>

I.e., with explicit emission of a copy of the pointed-to type following
the pointer itself.

When DFP is disabled, we instead get:

  <<< error >>>
  <<< error >>>
  <<< error >>>
  <<< error >>>
  <<< error >>>
  <<< error >>>

So, the number of nodes emitted during LTO write-out in the host/read-in
in the offload compiler do not match.

This patch restores the number of nodes emitted by creating
dfloatN_ptr_type_node as generic pointers rather than treating them as
flat error_type_nodes. I don't think there's an easy way of creating an
"error_type_node *", nor do I know if that would really be preferable.

Tested with offloading to NVPTX & bootstrapped. OK to apply?

Thanks,

Julian

ChangeLog

        gcc/
        * tree.c (build_common_tree_nodes): Use pointer type for
        dfloat32_ptr_type_node, dfloat64_ptr_type_node and
        dfloat128_ptr_type_node when decimal floating point support
        is disabled.

Comments

Joseph Myers Nov. 28, 2019, 3:04 p.m. | #1
On Thu, 28 Nov 2019, Julian Brown wrote:

> Unlike e.g. the _FloatN types, when decimal floating-point types are

> enabled, common tree nodes are created for each float type size (e.g.

> dfloat32_type_node) and also a pointer to each type is created

> (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node emits

> these like:


As far as I can tell, nothing actually uses those pointer nodes, or the 
corresponding BT_DFLOAT32_PTR etc. defined in builtin-types.def.  I don't 
know if they ever were used, or if they were just added by analogy to e.g. 
float_ptr_type_node.

So I'd suggest simply removing all references to those tree nodes and 
corresponding BT_*, from builtin-types.def, jit/jit-builtins.c (commented 
out), tree-core.h, tree.c, tree.h.  Hopefully that will solve the 
offloading problem.

-- 
Joseph S. Myers
joseph@codesourcery.com
Thomas Schwinge Nov. 28, 2019, 3:15 p.m. | #2
Hi Julian!

On 2019-11-28T14:24:02+0000, Julian Brown <julian@codesourcery.com> wrote:
> As mentioned in PR91985, offloading compilation is broken at present

> because of an issue with LTO streaming. With thanks to Joseph for

> hints, here's a solution.

>

> Unlike e.g. the _FloatN types, when decimal floating-point types are

> enabled, common tree nodes are created for each float type size (e.g.

> dfloat32_type_node) and also a pointer to each type is created

> (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node emits

> these like:

>

>   <float:32>     (dfloat32_type_node)

>   <float:64>     (dfloat64_type_node)

>   <float:128>    (dfloat128_type_node)

>   <float:32> *   (dfloat32_ptr_type_node)

>   <float:32>

>   <float:64> *   (dfloat64_ptr_type_node)

>   <float:64>

>   <float:128> *  (dfloat128_ptr_type_node)

>   <float:128>

>

> I.e., with explicit emission of a copy of the pointed-to type following

> the pointer itself.


I also do see that, but I fail to understand why that duplication: the
first '<float:32>' and the second one (after the '<float:32> *') are the
same node, or aren't they?

> When DFP is disabled, we instead get:

>

>   <<< error >>>

>   <<< error >>>

>   <<< error >>>

>   <<< error >>>

>   <<< error >>>

>   <<< error >>>


(With that expectedly being any 'NULL_TREE's converted to
'error_mark_node' in 'gcc/tree-streamer.c:record_common_node'.)

> So, the number of nodes emitted during LTO write-out in the host/read-in

> in the offload compiler do not match.


ACK.

> This patch restores the number of nodes emitted by creating

> dfloatN_ptr_type_node as generic pointers rather than treating them as

> flat error_type_nodes. I don't think there's an easy way of creating an

> "error_type_node *", nor do I know if that would really be preferable.

>

> Tested with offloading to NVPTX & bootstrapped. OK to apply?


> commit 17119773a8a45af098364b4faafe68f2e868479a

> Author: Julian Brown <julian@codesourcery.com>

> Date:   Wed Nov 27 18:41:56 2019 -0800

>

>     Fix decimal floating-point LTO streaming for offloading compilation

>     

>             gcc/

>             * tree.c (build_common_tree_nodes): Use pointer type for

>             dfloat32_ptr_type_node, dfloat64_ptr_type_node and

>             dfloat128_ptr_type_node when decimal floating point support is disabled.

>

> diff --git a/gcc/tree.c b/gcc/tree.c

> index 5ae250ee595..db3f225ea7f 100644

> --- a/gcc/tree.c

> +++ b/gcc/tree.c

> @@ -10354,6 +10354,15 @@ build_common_tree_nodes (bool signed_char)

>        layout_type (dfloat128_type_node);

>        dfloat128_ptr_type_node = build_pointer_type (dfloat128_type_node);

>      }

> +  else

> +    {

> +      /* These must be pointers else tree-streamer.c:record_common_node will emit

> +	 a different number of nodes depending on DFP availability, which breaks

> +	 offloading compilation.  */

> +      dfloat32_ptr_type_node = ptr_type_node;

> +      dfloat64_ptr_type_node = ptr_type_node;

> +      dfloat128_ptr_type_node = ptr_type_node;

> +    }

>  

>    complex_integer_type_node = build_complex_type (integer_type_node, true);

>    complex_float_type_node = build_complex_type (float_type_node, true);


(Maybe that's indeed better than my "hamfisted" patch.)  ;-)

But it still reads a bit like a workaround (explicitly setting only
'dfloat*_ptr_type_node' here, leaving the actual 'dfloat*_type_node'
untouched (and then later on implicitly converted to 'error_mark_node' as
mentioned).

I guess we need somebody with more experience to review this.


Grüße
 Thomas
Segher Boessenkool Nov. 28, 2019, 3:44 p.m. | #3
Hi Joseph,

On Thu, Nov 28, 2019 at 03:04:05PM +0000, Joseph Myers wrote:
> On Thu, 28 Nov 2019, Julian Brown wrote:

> > Unlike e.g. the _FloatN types, when decimal floating-point types are

> > enabled, common tree nodes are created for each float type size (e.g.

> > dfloat32_type_node) and also a pointer to each type is created

> > (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node emits

> > these like:

> 

> As far as I can tell, nothing actually uses those pointer nodes, or the 

> corresponding BT_DFLOAT32_PTR etc. defined in builtin-types.def.  I don't 

> know if they ever were used, or if they were just added by analogy to e.g. 

> float_ptr_type_node.

> 

> So I'd suggest simply removing all references to those tree nodes and 

> corresponding BT_*, from builtin-types.def, jit/jit-builtins.c (commented 

> out), tree-core.h, tree.c, tree.h.  Hopefully that will solve the 

> offloading problem.


So your patch caused at least three problems, none of them completely
worked out yet, none of them trivial.

Maybe this isn't such a good idea during stage 3.


Segher
Julian Brown Nov. 28, 2019, 6:21 p.m. | #4
On Thu, 28 Nov 2019 15:04:05 +0000
Joseph Myers <joseph@codesourcery.com> wrote:

> On Thu, 28 Nov 2019, Julian Brown wrote:

> 

> > Unlike e.g. the _FloatN types, when decimal floating-point types are

> > enabled, common tree nodes are created for each float type size

> > (e.g. dfloat32_type_node) and also a pointer to each type is created

> > (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node

> > emits these like:  

> 

> As far as I can tell, nothing actually uses those pointer nodes, or

> the corresponding BT_DFLOAT32_PTR etc. defined in builtin-types.def.

> I don't know if they ever were used, or if they were just added by

> analogy to e.g. float_ptr_type_node.

> 

> So I'd suggest simply removing all references to those tree nodes and 

> corresponding BT_*, from builtin-types.def, jit/jit-builtins.c

> (commented out), tree-core.h, tree.c, tree.h.  Hopefully that will

> solve the offloading problem.


Thanks for review. How about this (lightly retested so far), assuming
it passes full testing/bootstrap?

Thanks,

Julian

ChangeLog

        gcc/
        * builtin-types.def (BT_DFLOAT32_PTR, BT_DFLOAT64_PTR,
        BT_DFLOAT128_PTR) Remove.
        * tree-core.h (TI_DFLOAT32_PTR_TYPE, TI_DFLOAT64_PTR_TYPE,
        TI_DFLOAT128_PTR_TYPE): Remove.
        * tree.c (build_common_type_nodes): Remove dfloat32_ptr_type_node,
        dfloat64_ptr_type_node and dfloat128_ptr_type_node initialisation.
        * tree.h (dfloat32_ptr_type_node, dfloat64_ptr_type_node,
        dfloat128_ptr_type_node): Remove macros.

        gcc/jit/
        * jit-builtins.c (BT_DFLOAT32_PTR, BT_DFLOAT64_PTR, BT_DFLOAT128_PTR):
        Remove commented-out cases.
commit 80f69450724dbbf944a0d1e62e3ca6bdc3dd5a82
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed Nov 27 18:41:56 2019 -0800

    Remove unused decimal floating-point pointer types
    
            gcc/
            * builtin-types.def (BT_DFLOAT32_PTR, BT_DFLOAT64_PTR,
            BT_DFLOAT128_PTR) Remove.
            * tree-core.h (TI_DFLOAT32_PTR_TYPE, TI_DFLOAT64_PTR_TYPE,
            TI_DFLOAT128_PTR_TYPE): Remove.
            * tree.c (build_common_type_nodes): Remove dfloat32_ptr_type_node,
            dfloat64_ptr_type_node and dfloat128_ptr_type_node initialisation.
            * tree.h (dfloat32_ptr_type_node, dfloat64_ptr_type_node,
            dfloat128_ptr_type_node): Remove macros.
    
            gcc/jit/
            * jit-builtins.c (BT_DFLOAT32_PTR, BT_DFLOAT64_PTR, BT_DFLOAT128_PTR):
            Remove commented-out cases.

diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 800b751de6d..2611e88da60 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -145,15 +145,6 @@ DEF_PRIMITIVE_TYPE (BT_DFLOAT64, (dfloat64_type_node
 DEF_PRIMITIVE_TYPE (BT_DFLOAT128, (dfloat128_type_node
 				   ? dfloat128_type_node
 				   : error_mark_node))
-DEF_PRIMITIVE_TYPE (BT_DFLOAT32_PTR, (dfloat32_ptr_type_node
-				      ? dfloat32_ptr_type_node
-				      : error_mark_node))
-DEF_PRIMITIVE_TYPE (BT_DFLOAT64_PTR, (dfloat64_ptr_type_node
-				      ? dfloat64_ptr_type_node
-				      : error_mark_node))
-DEF_PRIMITIVE_TYPE (BT_DFLOAT128_PTR, (dfloat128_ptr_type_node
-				       ? dfloat128_ptr_type_node
-				       : error_mark_node))
 
 DEF_PRIMITIVE_TYPE (BT_VALIST_REF, va_list_ref_type_node)
 DEF_PRIMITIVE_TYPE (BT_VALIST_ARG, va_list_arg_type_node)
diff --git a/gcc/jit/jit-builtins.c b/gcc/jit/jit-builtins.c
index 850329c7b36..93d48c64c40 100644
--- a/gcc/jit/jit-builtins.c
+++ b/gcc/jit/jit-builtins.c
@@ -434,9 +434,6 @@ builtins_manager::make_primitive_type (enum jit_builtin_type type_id)
     // case BT_DFLOAT32:
     // case BT_DFLOAT64:
     // case BT_DFLOAT128:
-    // case BT_DFLOAT32_PTR:
-    // case BT_DFLOAT64_PTR:
-    // case BT_DFLOAT128_PTR:
     // case BT_VALIST_REF:
     // case BT_VALIST_ARG:
     // case BT_I1:
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 12e078882da..f76f68d835d 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -695,9 +695,6 @@ enum tree_index {
   TI_DFLOAT32_TYPE,
   TI_DFLOAT64_TYPE,
   TI_DFLOAT128_TYPE,
-  TI_DFLOAT32_PTR_TYPE,
-  TI_DFLOAT64_PTR_TYPE,
-  TI_DFLOAT128_PTR_TYPE,
 
   TI_VOID_LIST_NODE,
 
diff --git a/gcc/tree.c b/gcc/tree.c
index 5ae250ee595..789f0a00f41 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -10340,19 +10340,16 @@ build_common_tree_nodes (bool signed_char)
       TYPE_PRECISION (dfloat32_type_node) = DECIMAL32_TYPE_SIZE;
       SET_TYPE_MODE (dfloat32_type_node, SDmode);
       layout_type (dfloat32_type_node);
-      dfloat32_ptr_type_node = build_pointer_type (dfloat32_type_node);
 
       dfloat64_type_node = make_node (REAL_TYPE);
       TYPE_PRECISION (dfloat64_type_node) = DECIMAL64_TYPE_SIZE;
       SET_TYPE_MODE (dfloat64_type_node, DDmode);
       layout_type (dfloat64_type_node);
-      dfloat64_ptr_type_node = build_pointer_type (dfloat64_type_node);
 
       dfloat128_type_node = make_node (REAL_TYPE);
       TYPE_PRECISION (dfloat128_type_node) = DECIMAL128_TYPE_SIZE;
       SET_TYPE_MODE (dfloat128_type_node, TDmode);
       layout_type (dfloat128_type_node);
-      dfloat128_ptr_type_node = build_pointer_type (dfloat128_type_node);
     }
 
   complex_integer_type_node = build_complex_type (integer_type_node, true);
diff --git a/gcc/tree.h b/gcc/tree.h
index 60b6eae7b04..0f3cc5d7e5a 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -4098,9 +4098,6 @@ tree_strip_any_location_wrapper (tree exp)
 #define dfloat32_type_node              global_trees[TI_DFLOAT32_TYPE]
 #define dfloat64_type_node              global_trees[TI_DFLOAT64_TYPE]
 #define dfloat128_type_node             global_trees[TI_DFLOAT128_TYPE]
-#define dfloat32_ptr_type_node          global_trees[TI_DFLOAT32_PTR_TYPE]
-#define dfloat64_ptr_type_node          global_trees[TI_DFLOAT64_PTR_TYPE]
-#define dfloat128_ptr_type_node         global_trees[TI_DFLOAT128_PTR_TYPE]
 
 /* The fixed-point types.  */
 #define sat_short_fract_type_node       global_trees[TI_SAT_SFRACT_TYPE]
Joseph Myers Nov. 28, 2019, 6:25 p.m. | #5
On Thu, 28 Nov 2019, Julian Brown wrote:

> On Thu, 28 Nov 2019 15:04:05 +0000

> Joseph Myers <joseph@codesourcery.com> wrote:

> 

> > On Thu, 28 Nov 2019, Julian Brown wrote:

> > 

> > > Unlike e.g. the _FloatN types, when decimal floating-point types are

> > > enabled, common tree nodes are created for each float type size

> > > (e.g. dfloat32_type_node) and also a pointer to each type is created

> > > (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node

> > > emits these like:  

> > 

> > As far as I can tell, nothing actually uses those pointer nodes, or

> > the corresponding BT_DFLOAT32_PTR etc. defined in builtin-types.def.

> > I don't know if they ever were used, or if they were just added by

> > analogy to e.g. float_ptr_type_node.

> > 

> > So I'd suggest simply removing all references to those tree nodes and 

> > corresponding BT_*, from builtin-types.def, jit/jit-builtins.c

> > (commented out), tree-core.h, tree.c, tree.h.  Hopefully that will

> > solve the offloading problem.

> 

> Thanks for review. How about this (lightly retested so far), assuming

> it passes full testing/bootstrap?


This patch is OK.

-- 
Joseph S. Myers
joseph@codesourcery.com
Richard Biener Nov. 28, 2019, 7:40 p.m. | #6
On Thu, Nov 28, 2019 at 4:04 PM Joseph Myers <joseph@codesourcery.com> wrote:
>

> On Thu, 28 Nov 2019, Julian Brown wrote:

>

> > Unlike e.g. the _FloatN types, when decimal floating-point types are

> > enabled, common tree nodes are created for each float type size (e.g.

> > dfloat32_type_node) and also a pointer to each type is created

> > (e.g. dfloat32_ptr_type_node). tree-streamer.c:record_common_node emits

> > these like:

>

> As far as I can tell, nothing actually uses those pointer nodes, or the

> corresponding BT_DFLOAT32_PTR etc. defined in builtin-types.def.  I don't

> know if they ever were used, or if they were just added by analogy to e.g.

> float_ptr_type_node.

>

> So I'd suggest simply removing all references to those tree nodes and

> corresponding BT_*, from builtin-types.def, jit/jit-builtins.c (commented

> out), tree-core.h, tree.c, tree.h.  Hopefully that will solve the

> offloading problem.


Indeed that seems to be the case and would be my suggestion as well.
The issue with LTO streaming here is that pointers get streamed as two
things but the error-mark replacement as one, that causes the mismatches.

So please just remove those three global types.

Richard.

> --

> Joseph S. Myers

> joseph@codesourcery.com

Patch

commit 17119773a8a45af098364b4faafe68f2e868479a
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed Nov 27 18:41:56 2019 -0800

    Fix decimal floating-point LTO streaming for offloading compilation
    
            gcc/
            * tree.c (build_common_tree_nodes): Use pointer type for
            dfloat32_ptr_type_node, dfloat64_ptr_type_node and
            dfloat128_ptr_type_node when decimal floating point support is disabled.

diff --git a/gcc/tree.c b/gcc/tree.c
index 5ae250ee595..db3f225ea7f 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -10354,6 +10354,15 @@  build_common_tree_nodes (bool signed_char)
       layout_type (dfloat128_type_node);
       dfloat128_ptr_type_node = build_pointer_type (dfloat128_type_node);
     }
+  else
+    {
+      /* These must be pointers else tree-streamer.c:record_common_node will emit
+	 a different number of nodes depending on DFP availability, which breaks
+	 offloading compilation.  */
+      dfloat32_ptr_type_node = ptr_type_node;
+      dfloat64_ptr_type_node = ptr_type_node;
+      dfloat128_ptr_type_node = ptr_type_node;
+    }
 
   complex_integer_type_node = build_complex_type (integer_type_node, true);
   complex_float_type_node = build_complex_type (float_type_node, true);