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
