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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
