[SYCL] Support *global_[device|host]_space in static_address_cast

[aelovikov-intel]

When these address spaces are used with regular
sycl::detail::spirv::GenericCastToPtr they are turned into
unreachable. We don't have a SPIR-V intrinsic yet (or maybe we
shouldn't even have it, and will continue to rely on standard LLVM IR's
addrspacecast), so use C-style cast and rely on the translator/backend
to generate proper operation, similarly to sycl::detail::cast_AS.

[aelovikov-intel]

@gmlueck , @Pennycook , I'm not sure which extension should I modify for this, any suggestions?

[aelovikov-intel]

CI failures are caused by a recently merged PR, notified author/gatekeeper in #15407 (comment).

[Pennycook]

@gmlueck , @Pennycook , I'm not sure which extension should I modify for this, any suggestions?

I'm not sure if @gmlueck will agree, but part of me wonders if we can get away with not modifying anything. The ElementType* overloads defined for the address_cast functions already accept these pointers, and I think it's already clear which casts would fail...

If we have to change one of the specifications, I'd prefer to change the USM address one... But because that extension is "supported" it's probably easier to change the address cast one. We could just add a quick note that says something along the lines of:

"If an implementation supports sycl_ext_intel_usm_address_spaces, static_address_cast and dynamic_address_cast can be used to cast pointers in the ext_intel_global_device_space or ext_intel_global_host_space to global_space."

I don't think we want to define new shorthand aliases for the USM address spaces, or encourage their use.

[aelovikov-intel]

I don't think we want to define new shorthand aliases for the USM address spaces, or encourage their use

part of me wonders if we can get away with not modifying anything

I was confused by these aliases before and thought that the address cast extension was only supposed to work with these spaces. That does not seem to be the case, so I agree that we don't need any changes in the extension specification. (we might have a similar bug with constant AS though, I'll look at it separately).

[aelovikov-intel]

ping @dm-vodopyanov , @intel/llvm-reviewers-runtime

[aelovikov-intel]

@dm-vodopyanov , @intel/llvm-reviewers-runtime, ping x2.

[steffenlarsen]

Looks reasonable!

[steffenlarsen]

I think it's a bad habit of ours to have any other detail namespace than sycl::detail. Any details can be in sycl::detail, even if it's extension related. That said, not this PR's fault and I have added to those namespaces in the past. Just a bit of venting. 😉