yaxunl added a comment. In D77743#1972301 <https://reviews.llvm.org/D77743#1972301>, @hliao wrote:
> 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. Right. This will work. We don't need user to load from foo, because foo will resolve to kernel handle instead of stub function. 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
