[PATCH] D12246: [NVPTX] change threading intrinsics from noduplicate to convergent

[Owen Anderson via llvm-commits]

> On Aug 31, 2015, at 10:11 PM, Mehdi Amini <mehdi.amini at apple.com> wrote:
> 
> 
>> On Aug 31, 2015, at 10:01 PM, Owen Anderson <resistor at mac.com <mailto:resistor at mac.com>> wrote:
>> 
>> 
>>> On Aug 31, 2015, at 9:55 PM, Mehdi Amini <mehdi.amini at apple.com <mailto:mehdi.amini at apple.com>> wrote:
>>> 
>>>> Can you define what you mean by “predicates under which the original executed”?
>>> 
>>> I think I see what you mean: the syncthreads() call is made control dependent on b while it was not before. Your sentence made me think about something quite different.
>> 
>> There are multiple ways to view this problem, none of which are particularly native to LLVM.  Control dependence is one way.  Execution predicates (as seen in IRs like Sea of Nodes) is another, equivalent, way.
> 
> Sure, I was not trying to correct anything you wrote, just expressing that I finally understood what you meant :)
> I don’t know about Sea of Nodes and I associated the formulation you used to a totally different compiler concept than “execution predicates”. 
> 
> What about full unrolling?
> 
> for(int i = 0; i < 2; i++) {
>   __syncthreads();
> }
> 
> to
> 
> __syncthreads();
> __syncthreads();
> 
> This does not seem to preserve the "set of execution predicates”.
> (I’ll read about the "Sea of Nodes" model)

Yeah, this is where it gets murkier, and possibly where some variation of control dependence becomes more informative.  In some sense the dependence on i (or in the predicates world, the gating on the values of i) is trivial in that there’s no conditionality involved, so unrolling the loop does not introduce or remove conditionality.

—Owen

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150831/46163dfe/attachment.html>