[PATCH] D20039: [CUDA] Restrict init of local __shared__ variables to empty constructors only.

[Artem Belevich via cfe-commits]

tra added a comment.

In http://reviews.llvm.org/D20039#424067, @jlebar wrote:

> While I think this is 100% the right thing to do, I am worried about breaking existing targets.  Maybe we need an escape valve, at least until we get that sorted out?  Unless you're pretty confident this isn't happening / will be easy enough to fix.


While empty constructors for __shared__ variables are not unusual, I haven't seen any non-empty constructors in practice.
Considering that such constructor would crash clang until this patch, it would be hard to miss such cases.

Escape hatch would require to make clang to do something reasonable. 
Currently llvm crashes because we attempt to generate a load from a guard variable using atomic load instruction that's not supported by NVPTX backend. Even if it was supported, we'd also need to implement _cxa_guard_acquire/release.

We can disable thread-safe guard variants. This will be broken, too, because there will be many per-block instances of the variable, but only one global guard.

We can spend an effort to do the same thing as nvcc -- no guard of any kind, ctor/dtor on function entry/exit and compiler warning about guaranteed data race, but I don't see much point doing *that*.

IMO fixing the source code to initialize static variable explicitly would be the right thing to do.


http://reviews.llvm.org/D20039