[PATCH] D124382: [Clang] Recognize target address space in superset calculation

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia added a comment.

In D124382#3479773 <https://reviews.llvm.org/D124382#3479773>, @jchlanda wrote:

> In D124382#3472888 <https://reviews.llvm.org/D124382#3472888>, @Anastasia wrote:
>
>> 
>
>
>
>> Can you provide an example of where it could be useful? Note that I feel that
>> such functionality could be implemented on top of full implementation of
>> target specific address space proposed in https://reviews.llvm.org/D62574.
>
> The use case we had was when calling target builtin (that specifies address
> space) from within OpenCL C. Currently this errors out, similarly, the explicit
> type cast to address space yields an error
> (`(__attribute__((address_space(3))) int *)&l_woof` in the example below). This
> is important for libclc, which is implemented in OpenCL C and deals directly
> with target builtins.
>
>   __kernel void woof() {
>     __local int l_woof;
>     __nvvm_cp_async_mbarrier_arrive_shared(&l_woof);
>   }
>
> I wasn't aware of that patch, sorry, I've not had a close look yet, but it
> seems worryingly dated.

Ok, I think the current behavior of builtins is to work with any address space. The way it had worked so far is since the builtins are only intended to be used in toolchains (instead of arbitrary code), the toolchain developers were responsible for making sure that the address spaces are used adequately in those builtins. However the question of extending the clang builtin functions with the notion of language address spaces has popped up before. And I think we could add this feature in a very light way for example by reserving the numbers from the clang `LangAS` enum to be used with the language address spaces in the prototypes of builtins. Although we could think of more elegant alternatives too. My understanding is there was never a strong enough case to add this functionality.

Just to understand further - do you need the builtins to have a specific address space prototype? And do you need this to improve error handling in libclc code base?

I imagine you could also create some sort of the wrapper functions around the building with the right address spaces, i.e. something like

  void __libclc_builtin1(local int* p){
     __builtin1(p);
  }

So the prototype in Clang for `__builtin1` would still be permissive wrt address space of the pointer but as you only use `__libclc_builtin1` in the codebase, you can ensure the correct uses. While this is how this problem has been worked around I think extending Clang builtins definitions might be inevitable to avoid forcing toolchains to create wrapper functions. However if we are aiming for this goal, I think more targeted solutions would make more sense instead of solving this problem indirectly by allowing conversions between target and language address spaces as this for example won’t work for builtins that are shared between targets.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D124382/new/

https://reviews.llvm.org/D124382