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

[Bjarke Roune via llvm-commits]

On Fri, Aug 28, 2015 at 5:15 PM, escha <escha at apple.com> wrote:
>
> [...] I'm still concerned about loop unrolling in a case such as this:

>
> >  for (int i = 0; i < *bound; ++i) {
> >    if (i == 0)
> >      __syncthreads();
> >  }
> >
> > This input program is valid as long as *bound > 0 has the same value
> across the block. Here loop-unrolling by a factor of 2 will separate off
> the first iteration of the loop into a duplicate body for the case where
> *bound is odd. I checked with an example loop that's similar but that
> doesn't use __syncthreads() and LLVM does do unrolling by a factor of 2 in
> this way. If whether *bound is odd is divergent, then only part of the warp
> would execute the __syncthreads() in the duplicate odd-case unrolled loop
> body.
>
> This would be a violation of convergent semantics, IMO. Anything that
> causes a convergent instruction that was uniform to become not-uniform is,
> I think, duplication or not.
>
> —escha


I agree. What my new example shows is that loop unrolling should take the
convergent attribute into account, but it doesn't currently do that AFAICT.
Sinking::SinkInstruction is the only place that I've found where the
convergent attributed is considered for IR right now.

Loop unswitching also needs to take the convergent attribute into account.
Take this example:

  bool b = *ptr;
  for (int i = 0; i < 100; ++i) {
    if (b) {
      // do something
    }
    __syncthreads();
  }

Loop unswitching could decide to unswitch on b even if b is divergent,
which would duplicate the `__syncthreads()`. That's normally a performance
problem, but with convergent/__syncthreads() it's a correctness problem.

The langref doesn't give semantics for when instructions marked convergent
can be duplicated, only for when they can be moved. Owen, do you have some
semantics for this in mind?

Bjarke
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150828/4ee3e6ed/attachment.html>