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

[Michael Liao via Phabricator via cfe-commits]

hliao added a comment.

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.


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

https://reviews.llvm.org/D77743