[PATCH] D85603: IR: Add convergence control operand bundle and intrinsics

[Sameer Sahasrabuddhe via Phabricator via llvm-commits]

sameerds added a comment.

In D85603#2370240 <https://reviews.llvm.org/D85603#2370240>, @vgrover99 wrote:

> I believe what is described here about convergent, as best I can understand it, is the semantics of __syncthreads in CUDA. This semantics is the same for <sm70 and sm70+. Not clear whether what is described here is a "textually aligned" semantics or unaligned. __syncthreads is aligned, meaning that all threads in the threadblock must wait on the same lexical syncthreads().

Textual alignment is a good way to examine this spec with respect to CUDA. The notion of dynamic instances is textually aligned according to basic rule 2:

> 2. Executions of different instructions always occur in different dynamic instances. For this and other rules in this document, instructions of the same type at different points in the program are considered to be different instructions.

This correctly covers `__syncthreads()`, and is a bit conservative about builtins that take a mask like `__syncwarp()`. In @jlebar's example, each call to `__syncwarp()` is a separate dynamic instance, although CUDA actually treats them as a single synchronization point.

  // Warning, does not seem to work on sm75
  if (cond)
    __syncwarp(FULL_MASK);
  else
    __syncwarp(FULL_MASK);

In general, the formal rules work correctly for this too: hoisting and sinking is disallowed without additional information. So the proposal is compatible with the new CUDA semantics for these builtins. These builtins do need convergence control: sinking such a call across a condition should be forbidden by default, since we can no longer guarantee that every thread in the mask still makes a matching call. Of course, specific optimizations that can recompute the mask can over-ride this restriction.

  foo = __shfl_sync(mask, ...);
  if (condition) {
    // cannot sink foo here if condition is divergent
    sole_use_of_foo();
  }


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D85603/new/

https://reviews.llvm.org/D85603