[llvm-commits] DAG type Legalizer bug?

Thu Apr 19 15:23:06 PDT 2012

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,v3f32,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>
> >>>>
> >>>
> >>>
> >>>
> >>
> >
> >
> 