| Message ID | 20191128142402.2949931d@squid.athome |
|---|---|
| State | New |
| Headers | show |
| Series | Fix decimal floating-point LTO streaming for offloading compilation | expand |
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.
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
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
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.
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.
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
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);
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.