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

Mehdi Amini via llvm-commits llvm-commits at lists.llvm.org
Mon Aug 31 22:11:20 PDT 2015


> On Aug 31, 2015, at 10:01 PM, Owen Anderson <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)


— 
Mehdi


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


More information about the llvm-commits mailing list