arsenm wrote: > > I'd rather just fix clang. This is a workaround that shouldn't be > > necessary, and seems to spread the knowledge of the dodgy > > !__opencl_c_generic_address_space address space handling case to a new place > > By fixing clang you mean the cast `(void [[clang::address_space(0)]] *)((void > [[clang::opencl_generic]] *)ptr)` should compile in all language modes?
Well it certainly shouldn't crash...the exact semantics I'm less sure of. The behavior of the numbered address spaces here is emergent behavior, I don't think any of this has any intentional. Proper error would be an improvement over the status quo. If there's an easy backdoor to get the addrspacecast emitted to generic if you try to use these in OpenCL 1.x that would probably be the nicest option. I think the ultimate fix would be to stop treating the default address space as private in all language modes. That is, OpenCL 1.x would be better implemented by having an implicit clang::opencl_private attribute applied to all pointers, with opencl_generic as the default address space, rather than the language mode changing this fundamental property. > Or change the `__builtin_amdgcn_is_(private|shared)` instrinsic to accept > pointers in the default (language) address space instead of the target > address space 0? These builtins should only accept language address spaces that will resolve to the 0 IR address space. It would be wrong to take the OpenCL 1.x interpretation of "default" as a private pointer and accept that here. That is, some notion of generic address space needs to exist to call these. Whether there should be a backdoor in OpenCL 1.x to insert the cast, I'm not sure. There's no real association between the language address space and the numbered address spaces and it's debatable if there should be. From what I've gathered, the address space attribute is for "user" address spaces, so one possible interpretation is we could renumber those to some other IR address spaces and not treat them as aliases. However in real code there is already widespread abuse of numbered address spaces to get at these, which bypasses all the useful semantic checks. It might be worthwhile to teach clang to treat the target numbers as aliases for the proper language address_spaces https://github.com/llvm/llvm-project/pull/152314 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
