[PATCH] D77743: [HIP] Emit symbols with kernel name in host binary

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D77743#1972301 <https://reviews.llvm.org/D77743#1972301>, @hliao wrote:

> In D77743#1972298 <https://reviews.llvm.org/D77743#1972298>, @yaxunl wrote:
>
> > In D77743#1972292 <https://reviews.llvm.org/D77743#1972292>, @hliao wrote:
> >
> > > In addition, we may also need to extend the registration to set up the mapping from that global variable to the host side stub function. `hipKernelLaunch` (implemented as a function call instead of the kernel launch syntax) to call into that stub function to prepare the arguments.
> >
> >
> > hipKernelLaunch does not call the stub function. The stub function calls hipKernelLaunch. Therefore user/runtime does not need to know about stub function to launch a kernel.
>
>
> Since the code using hipKernelLuanch may be compiled by other compilers, we cannot force reinterpreting the use of that symbol by loading value from the symbol. For code like this
>
>   __global__ void foo();
>  
>   hipKernelLaunch(foo, ...)
>
>
> If it's compiled by other compiler, `foo` refers to the value of that symbol, i.e. a constant, instead of the value loading from that symbol. They are different.


Right. This will work. We don't need user to load from foo, because foo will resolve to kernel handle instead of stub function.


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

https://reviews.llvm.org/D77743