…V path

From SYCL 2020 specification:

> The sycl::atomic_ref class also has a template parameter AddressSpace,
> which allows the application to make an assertion about the address
> space of the object of type T that it references. The default value
> for this parameter is access::address_space::generic_space, which
> indicates that the object could be in either the global or local
> address spaces. If the application knows the address space, it can set
> this template parameter to either access::address_space::global_space
> or access::address_space::local_space as an assertion to the
> implementation. Specifying the address space via this template
> parameter may allow the implementation to perform certain
> optimizations. Specifying an address space that does not match the
> object’s actual address space results in undefined behavior

We use `ext::oneapi::experimental::static_address_cast` to do that. It's
not implemented for CUDA/HIP yet, that path continues using
`sycl::address_space_cast` that performs runtime checks:

> An implementation must return nullptr if the run-time value of pointer
> is not compatible with Space, and must issue a compiletime diagnostic
> if the deduced address space for pointer is not compatible with Space.

SPIRV operations are defined such that `OpGenericCastToPtr` and
`OpGenericCastToPtrExplicit` cannot be used when target `Storage Class`
is `Generic`, yet we were generating such code. This PR fixes that.