[llvm-commits] DAG type Legalizer bug?

Thu Apr 19 15:23:52 PDT 2012

Sorry, the test case is simplified, which is why the assert looks different.

> -----Original Message-----
> From: Villmow, Micah
> Sent: Thursday, April 19, 2012 3:23 PM
> To: 'Dan Gohman'; Guo, Xiaoyi
> Cc: Evan Cheng; llvm-commits at cs.uiuc.edu (llvm-commits at cs.uiuc.edu)
> Subject: RE: [llvm-commits] DAG type Legalizer bug?
> 
> Here is the original test case that caused the issue:
> typedef struct
> {
>   // Offset between the pixel and the sampled position
>   float3 offsetX;
>   uint pixelScramble
> } clTask;
> 
> __kernel __attribute__((vec_type_hint(float4))) void
> kernel__CreateCameraRays( __global clTask* tasks, __global clTask*
> tasks_out) {
>   size_t gid = get_global_id(0);
>   tasks_out[gid] = tasks[gid];
> }
> 
> It was hitting this assert in vector types:
> aoc ray_crash.cl -g -cl-opt-disable
> WidenVectorResult #8: 0x4084ec0:
> i32,i32,i32,f32,f32,i32,i32,i32,v3f32,v3f32,i32
> ,i32,f32,f32,v3f32,v3f32,f32,f32,i32,i32,f32,f32,f32,i32,v3f32,v3f32,v3f
> 32,v3f32
> ,i32,i32,i32,i32,v3f32,v3f32,i32,i32,i32,i32 = merge_values 0x407e780,
> 0x407e3c8
> , 0x4085278, 0x407ec48, 0x4075b00, 0x4085c98, 0x4076408, 0x4080880,
> 0x407e560,
> 0
> x4080908, 0x40772f0, 0x407e918, 0x4084a80, 0x4085388, 0x4085968,
> 0x408df88,
> 0x40
> 80e58, 0x4085c10, 0x408e450, 0x40802a8, 0x4076eb0, 0x4077268, 0x4085a78,
> 0x4075c 98, 0x408cd70, 0x4080990, 0x4084c18, 0x4085eb8, 0x4077400,
> 0x4080ee0, 0x4084860,  0x4086380, 0x408e4d8, 0x4077d10, 0x4076da0,
> 0x407e890, 0x407eab0, 0x40859f0 [OR D=732] [ID=0]
> dbg:OCLF5E4.tmp.cl:1806:42
> 
> Do not know how to widen the result of this operator!
> UNREACHABLE executed at ..\..\..\LegalizeVectorTypes.cpp:1244!
> 
> So as long as this test case still works.
> 
> I was not able to reproduce with the X86 backend however.
> 
> > -----Original Message-----
> > From: Dan Gohman [mailto:gohman at apple.com]
> > Sent: Thursday, April 19, 2012 3:19 PM
> > To: Guo, Xiaoyi
> > Cc: Villmow, Micah; Evan Cheng; llvm-commits at cs.uiuc.edu (llvm-
> > commits at cs.uiuc.edu)
> > Subject: Re: [llvm-commits] DAG type Legalizer bug?
> >
> > I'm not really familiar enough with the history of this code to have a
> > strong opinion on how it ought to work.
> > If Micah thinks your fix is fine, and if it doesn't break other stuff,
> > then it's probably fine.
> >
> > Dan
> >
> > On Apr 17, 2012, at 1:42 PM, "Guo, Xiaoyi" <Xiaoyi.Guo at amd.com> wrote:
> >
> > > Thanks Micah for confirming it.
> > >
> > > Dan,
> > >
> > > My fix was exactly to restore the old behavior by returning the
> > operand, not the value itself. If you agree with my fix, I'll fix the
> > comment also.
> > >
> > > Xiaoyi
> > >
> > > -----Original Message-----
> > > From: Villmow, Micah
> > > Sent: Tuesday, April 17, 2012 10:11 AM
> > > To: Guo, Xiaoyi; Dan Gohman
> > > Cc: Evan Cheng; llvm-commits at cs.uiuc.edu (llvm-commits at cs.uiuc.edu)
> > > Subject: RE: [llvm-commits] DAG type Legalizer bug?
> > >
> > > Here is the original patch:
> > > http://lists.cs.uiuc.edu/pipermail/llvm-commits/Week-of-Mon-20110822
> > > /1
> > > 27062.html
> > > Here is the second patch:
> > > http://lists.cs.uiuc.edu/pipermail/llvm-commits/Week-of-Mon-20110829
> > > /1
> > > 27142.html
> > >
> > > The fix was for this bug:
> > > http://llvm.org/bugs/show_bug.cgi?id=10736
> > >
> > > DecomposeMERGE_VALUES was written to return the operand that needs
> > > to
> > be modified, not the value. It looks like that what replaced it does
> > not do the same thing.
> > >
> > > This was the old behavior:
> > > +  // DecomposeMERGE_VALUES  takes a SDNode and returns the first
> > > + // operand that needs to be modified.
> > > +  // All other nodes are legalized, whether they are legal or not.
> > > +  // The resulting SDValue needs to be modified to make it legal.
> > >
> > > Here is the new behavior:
> > > /// DisintegrateMERGE_VALUES - Replace each result of the given
> > MERGE_VALUES /// node with the corresponding input operand, except for
> > the result 'ResNo', /// which is returned.
> > >
> > >> -----Original Message-----
> > >> From: Guo, Xiaoyi
> > >> Sent: Monday, April 16, 2012 8:06 PM
> > >> To: Dan Gohman
> > >> Cc: Evan Cheng; llvm-commits at cs.uiuc.edu
> > >> (llvm-commits at cs.uiuc.edu); Villmow, Micah
> > >> Subject: RE: [llvm-commits] DAG type Legalizer bug?
> > >>
> > >> The bug wasn't there at r138877, but was introduced later in
> > >> r140373 when DecomposeMERGE_VALUES() was replaced by
> > DisintegrateMERGE_VALUES().
> > >>
> > >> Since Micah is the original owner of this area, I'll ask him to
> > confirm.
> > >>
> > >> Thanks,
> > >> Xiaoyi
> > >>
> > >> -----Original Message-----
> > >> From: Dan Gohman [mailto:gohman at apple.com]
> > >> Sent: Monday, April 16, 2012 4:50 PM
> > >> To: Guo, Xiaoyi
> > >> Cc: Evan Cheng; llvm-commits at cs.uiuc.edu
> > >> (llvm-commits at cs.uiuc.edu); Villmow, Micah
> > >> Subject: Re: [llvm-commits] DAG type Legalizer bug?
> > >>
> > >> I get that there's a bug somewhere, but the question is whether the
> > >> bug is in DisintegrateMERGE_VALUES or SplitRes_MERGE_VALUES. I
> > >> don't know if there's a special reason why DisintegrateMERGE_VALUES
> > >> was written the way it is, but it at least doesn't seem accidental.
> > >>
> > >> svn blame attributes the code in SplitRes_MERGE_VALUES r138877,
> > >> which was a patch by Micah. Perhaps he knows more about this.
> > >>
> > >> Dan
> > >>
> > >> On Apr 16, 2012, at 1:31 PM, "Guo, Xiaoyi" <Xiaoyi.Guo at amd.com>
> > wrote:
> > >>
> > >>> Yes, I agree the code does what the comment says it does. But
> > >>> still,
> > >> as I explained in my original email below, there is a bug if you
> > >> look at how the returned value is used. To prove my point from
> > >> another point of view, if you look at the svn history of this
> > >> change, the comment for the revision doesn't show the intention to
> > >> change the behavior, however, if you look at the change, it did
> > >> change the behavior by returning the N node itself instead of
> > >> returning the
> > operand.
> > >>>
> > >>> If you agree with my assessment and my fix, I will update the
> > >>> comment
> > >> of the routine to agree with the fix.
> > >>>
> > >>> Xiaoyi
> > >>>
> > >>> -----Original Message-----
> > >>> From: Dan Gohman [mailto:gohman at apple.com]
> > >>> Sent: Friday, April 13, 2012 5:04 PM
> > >>> To: Evan Cheng
> > >>> Cc: Guo, Xiaoyi; llvm-commits at cs.uiuc.edu
> > >>> Subject: Re: [llvm-commits] DAG type Legalizer bug?
> > >>>
> > >>> I'm not deeply familiar with that code, but the current code in
> > >> DisintegrateMERGE_VALUES does precisely what its doxygen comment
> > >> says it does, which suggest that it's not written that way by
> accident.
> > >>>
> > >>> Dan
> > >>>
> > >>> On Apr 13, 2012, at 12:05 PM, Evan Cheng <evan.cheng at apple.com>
> > wrote:
> > >>>
> > >>>> Dan, can you review the patch?
> > >>>>
> > >>>> Thanks,
> > >>>>
> > >>>> Evan
> > >>>>
> > >>>> On Apr 13, 2012, at 10:58 AM, "Guo, Xiaoyi" <Xiaoyi.Guo at amd.com>
> > >> wrote:
> > >>>>
> > >>>>> No, this has not been reviewed. I attached the patched again.
> > >>>>>
> > >>>>> Thanks,
> > >>>>> Xiaoyi
> > >>>>>
> > >>>>> From: Evan Cheng [mailto:evan.cheng at apple.com]
> > >>>>> Sent: Thursday, April 12, 2012 10:27 PM
> > >>>>> To: Guo, Xiaoyi
> > >>>>> Cc: llvm-commits at cs.uiuc.edu
> > >>>>> Subject: Re: [llvm-commits] DAG type Legalizer bug?
> > >>>>>
> > >>>>> Has this been reviewed? Can you post the patch again?
> > >>>>>
> > >>>>> Evan
> > >>>>>
> > >>>>> On Apr 5, 2012, at 2:23 PM, "Guo, Xiaoyi" <Xiaoyi.Guo at amd.com>
> > >> wrote:
> > >>>>>
> > >>>>> Ping?
> > >>>>>
> > >>>>> From: llvm-commits-bounces at cs.uiuc.edu
> > >>>>> [mailto:llvm-commits-bounces at cs.uiuc.edu] On Behalf OfGuo,
> > >>>>> Xiaoyi
> > >>>>> Sent: Tuesday, April 03, 2012 11:39 AM
> > >>>>> To: llvm-commits at cs.uiuc.edu
> > >>>>> Subject: [llvm-commits] FW: DAG type Legalizer bug?
> > >>>>>
> > >>>>> Please review the attached patch which fixes a bug as described
> > >> below. Our test case fails on amdil backend. I failed to create a
> > >> test case with one of the backends built-in to llvm. So I couldn't
> > >> add a test case to the unit test suite.
> > >>>>>
> > >>>>> Thanks,
> > >>>>> Xiaoyi
> > >>>>>
> > >>>>> From: llvmdev-bounces at cs.uiuc.edu
> > >>>>> [mailto:llvmdev-bounces at cs.uiuc.edu] On Behalf Of Guo, Xiaoyi
> > >>>>> Sent: Friday, March 23, 2012 8:11 PM
> > >>>>> To: LLVM Developers Mailing List
> > >>>>> Subject: [LLVMdev] DAG type Legalizer bug?
> > >>>>>
> > >>>>> The following looks like a bug in the legalizer to me.
> > >>>>>
> > >>>>> DAGTypeLegalizer::SplitRes_MERGE_VALUES(SDNode*N, unsigned
> > >>>>> ResNo, SDValue& Lo, SDValue& Hi) {  SDValue Op =
> > >>>>> DisintegrateMERGE_VALUES(N, ResNo);  GetSplitOp(Op, Lo, Hi); }
> > >>>>>
> > >>>>> DisintegrateMERGE_VALUE() returns SDValue(N, ResNo), where N is
> > >>>>> the
> > >> MERGE_VALUE node itself.
> > >>>>> Then GetSplitOp() tries to retrieve split result for N from the
> > >> SplitVectors cache and hit assert because split result is for N is
> > >> not in the cache yet.
> > >>>>>
> > >>>>> Seems to me that DisintegrateMERGE_VALUES() should return the
> > >> corresponding operand for the given ResNo, not the defined value.
> > >>>>>
> > >>>>> Please confirm if it's a bug, or if I'm missing something.
> > >>>>>
> > >>>>> Thanks,
> > >>>>> Xiaoyi
> > >>>>> _______________________________________________
> > >>>>> llvm-commits mailing list
> > >>>>> llvm-commits at cs.uiuc.edu
> > >>>>> http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits
> > >>>>> <LegalizeTypes.patch>
> > >>>>
> > >>>
> > >>>
> > >>>
> > >>
> > >
> > >
> >