[clang] [clang-tools-extra] [Clang] Use TargetInfo when deciding is an address space is compatible (PR #115777)

[Joseph Huber via cfe-commits]

jhuber6 wrote:

> > For AMDGPU and NVPTX, pointers are allowed to drop their address spaces transparently. So, `void AS(3) *` can convert to `void *`. Languages like CUDA and SYCL and HIP use these 'language' address spaces to describe exactly this, but my assertion is that it's more of a function of the target. I'm sure @arsenm could explain the underlying implementation more deeply.
> 
> I think this view is somewhat problematic, as it mixes target concepts and language concepts in a dangerous way. Neither CUDA, nor HIP nor SYCL (AFAICT) do what is described, all of those are AS agnostic _languages_. Unlike e.g. OpenCL, a `__shared__` `int` is just an `int`, whereas the way in which we handle ASes in Clang / on the language level is typeful i.e. a `__private` `int` is not an `int` in OpenCL. An explicit pointer to a numbered AS generally does not blindly convert to a vanilla pointer in HIP or CUDA, observe e.g.: https://gcc.godbolt.org/z/cY4aa9cx8. So no, HIP and CUDA don't actually allow this behaviour per se. They chose different linguistic defaults. Furthermore, it's tricky to uplift target semantics into linguistic ones, for a language that is supposed to be generic.
>

I could see only applying these rules to the target specific (i.e. non-language AS's) so `AS(3) *` would convert to `*` but not `__shared__ *`. The above example doesn't work because we consider `cuda_shared` and `AS(3)` distinct even though they lower to the same thing for the target. 

> I don't necessarily mind this in this form, as an opt-in "here-be-dragons" mechanism which gets one into rather unspecified non-standard language dialect/extension territory. I do object to the very relaxed take that it is sound and just how things should work because some targets allow it. We should probably not try to sneakily back-legislate what standard C/C++ are.

I would greatly prefer this not to be opt in, but I could at least live with it if that's the ultimate decision. The take here is that it's only allows *on* targets that support it. We're not unilaterally defining that address spaces can be converted, it's just stating that we should use the target information when defining which casts are legal. I could update the logic in the `isAddrspaceSuperSetOf` function to be more explicit. For example ignoring language AS's and returning false on ones outside of our expected range (i.e. AS(999) is illegal).

My generous reading of the C statements was that casting legality was dependent on the target, ergo I figured this was a fine solution. But I'm not a language lawyer so I could be incorrect.

https://github.com/llvm/llvm-project/pull/115777