Synchronization in union-split code paths

[maleadt]

@tkf recently noticed / ran into a deadlock due to synchronization from branches that ended up divergent because of the compiler performing union splitting. To remedy this, we could use barrier.sync instead of barrier.sync.aligned (aka. bar.sync aka. syncthreads):

Instruction barrier has optional .aligned modifier. When specified, it indicates that all threads in CTA will execute the same barrier instruction. In conditionally executed code, an aligned barrier instruction should only be used if it is known that all threads in CTA evaluate the condition identically, otherwise behavior is undefined.

I'm not sure it's safe to just have syncthreads use an unaligned barrier, because we might then be inadvertently matching unrelated barriers:

if (foo)
  sync_threads() # thread 1 waits here
sync_threads()   # thread 2 waits here

Maybe that only matters for already broken code though, since thread 1 would here just get stuck at the second barrier?

Also interesting: at the PTX level, barriers are identified by a 'name' (int in 0:16) and a thread mask, so we could expose a more fine-grained sync for use in divergent branches. Maybe we could use this to differentiate barriers, but then we'd need some sort of lexical information to assign a barrier name, and I'd rather not move to a @sync_threads macro.

cc @vchuravy

[maleadt]

I guess its not semantically equivalent. This would deadlock with barrier.sync.aligned (disregarding compiler optimizations here):

x = rand()
if x < 0.5
  sync_threads()
else
  sync_threads()

But wouldn't without. Again, a case of code that was bad in the first place though. I wonder of semantics of legit code could change.

[vchuravy]

I wonder if there is an optimization to turn barrier.sync into barrier.sync.aligned if LLVM can prove it to be convergent (I recall something like this being discussed for the AMDGPU backend). In general I do think emitting the unaligned variant is safer.

Alternativly we could sink the sync_threads into the sucessor for union splitting.

[maleadt]

Alternativly we could sink the sync_threads into the sucessor for union splitting.

That's not possible in generality, right? The union-split callee could be an arbitrarily complex call (say, a mapreduce invocation). Of course, in that scenario we could end up further diverging by dispatching to different methods, which theoretically could be synchronizing differently...

[tkf]

Alternativly we could sink the sync_threads into the sucessor for union splitting.

That's not possible in generality, right? The union-split callee could be an arbitrarily complex call (say, a mapreduce invocation)

Yeah, I agree. It sounds difficult when sync_threads happens inside a loop (as would be done in mapreduce implementations) and the loop is inside of a function (so that union split happens outside the loop).

---

@vchuravy told me that he just found this: An abstract interpretation for SPMD divergence on reducible control flow graphs https://dl.acm.org/doi/abs/10.1145/3434312

---

With the function boundaries, I think we have two somewhat different problems: union splitting and method dispatch. I was focusing on union splitting because I felt it was very "invisible."

(I also wonder if this is a very Julia-specific problem. Are there languages that (1) can be compiled to GPU, (2) can do multi-argument function overloading, and (3) specialize for all arguments? e.g., I think C++ satisfies 1 and 2 but not 3.)

With llvm.nvvm.barrier.sync as suggested in #798, IIUC, it'd also "solve" the problem with method dispatch; i.e., if a function has multiple methods with different implementations that can synchronize each other using sync_threads. I don't know if this is a solution, though. It sounds dangerous. I wonder if you can use llvm.nvvm.barrier.sync only if the source Julia function call stack is using identical methods. Not sure if it's worth the hustle, though. Naively thinking, #798 looks like it does not change the behavior of already valid programs and makes some more programs valid.

---

BTW, thanks a lot for finding llvm.nvvm.barrier.sync! This is very exciting for me, since the only way to ask the compiler to union-split at return position, AFAICT, is to use the continuation-passing style which introduces a lot of function boundaries. I think llvm.nvvm.barrier.sync lets me execute a lot more kinds of state machines generated by transducers on GPU.