[PATCH] D147719: [OpenMP] Replace HeapToShared's initial value with `poison`

[Artem Belevich via Phabricator via llvm-commits]

tra added a comment.

Can you elaborate on why this would be the right thing to do?

We generally do want to treat shared memory loads/stores as `volatile` and block optimizations on them, as shared memory is often modified from other threads. In that sense neither poison nor undef appear to reflect that. I guess we may need to `freeze` all values loaded from shared memory.

E.g. if we have code like this, use of 'x' is fine, even though to LLVM it would look as if we always read uninitialized value from x.

  __shared__ int x;
  __global__ void kernel() {
    if (threadIdx.x == 0) {
       x = compute();
    }
    syncthreads();
    if (threadIdx.x != 0) {
      if (x != 42)  // LLVM should not be allowed to assume any specific value for x, and must do the actual comparison here.
        do_something(x);
    }
  }


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D147719