[CUDA] clang calls __cxa_guard* on the GPU side

[Artem-B]

Dynamic initializers are not allowed on the GPU.

CUDA allows using empty initializers, but when we actually have one, it goes through the regular codegen and ens up being guarded by __cxa_guard_* calls, if the empty initializer is not eliminated by the optimizations.

E.g. https://godbolt.org/z/P67j7M81c

This also exposes inconsistent behavior vs nvcc.

When initializer is not empty, clang does allow for static local variables (and that does need a guard), but NVCC rejects it, presumably because it would require guarding it.

I think we need to
a) make sure we do not emit thread guards for the static initializers
b) disallow non-empty initializers for static local variables (that would include local __shared__ variables that are implicitly static).