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



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to