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

[Artem Belevich via Phabricator via cfe-commits]

tra accepted this revision.
tra added a comment.
This revision is now accepted and ready to land.

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.


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

https://reviews.llvm.org/D77743