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

[Michael Liao via Phabricator via cfe-commits]

hliao requested changes to this revision.
hliao added a comment.
This revision now requires changes to proceed.

In D77743#1970304 <https://reviews.llvm.org/D77743#1970304>, @tra wrote:

> In D77743#1970163 <https://reviews.llvm.org/D77743#1970163>, @yaxunl wrote:
>
> > The kernel handle is a variable. Even if it has the same name as kernel, it is OK for the debugger since the debugger does not put break point on a variable.
>
>
> The patch appears to apply only to generated kernels. What happens when we take address of the kernel directly?
>
>   a.hip: 
>   __global__ void kernel() {}
>  
>   auto kernel_ref() {
>     return kernel;
>   }
>  
>   b.hip:
>   extern __global__ void kernel(); // access the handle var
>   something kernel_ref(); // returns the stub pointer?
>  
>   void f() {
>     auto x = kernel_ref();
>     auto y = kernel(); 
>     hipLaunchKernel(x,...); // x is the stub pointer. 
>     hipLaunchKernel(y,...);
>   }
>
>
> Will `x` and `y` contain the same value? For CUDA the answer would be yes as they both would contain the address of the host-side stub with the kernel's name.
>  In this case external reference will point to the handle variable, but I'm not sure what would kernel_ref() return. 
>  My guess is that it will be the stub address, which may be a problem.  I may be wrong. It would be good to add a test to verify that we always get consistent results when we're referencing the kernel.


That's a good question. That introduces the ambiguity on the values of the same symbol (from the programmer point of view). To ensure we won't there's no ambiguity, we should always use that *alias* global variable for `__global__` function on the host side as it will be used in the runtime API to query the device-side function.


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

https://reviews.llvm.org/D77743