tra added a comment. In D86376#2239391 <https://reviews.llvm.org/D86376#2239391>, @yaxunl wrote:
> For example, in HIP program, there is a kernel `void foo(int*)`. If a C++ > program wants to launch it, the desirable way is > > void foo(int*); > hipLaunchKernel(foo, grids, blocks, args, shmem, stream); > > Due to the prefixed kernel stub name, currently the users have to use > > void __device_stub_foo(int*); > hipLaunchKernel(__device_stub_foo, grids, blocks, args, shmem, stream); Ah. That *is* painful. Perhaps we can have the cake and eat it here and do something like this: Do generate a variable with the kernel name and use it for hipKernelLaunch(), but also keep the stub and call it for `<<<>>>` launches, only instead of using the stub itself registered as the GPU-side kernel identifier, use the variable. This way, __device_stub_<kernel> will show up in the stack trace (no debuggability regression), but direct calls to hipLaunchKenrel can use unprefixed kernel name. WDYT? CHANGES SINCE LAST ACTION https://reviews.llvm.org/D86376/new/ https://reviews.llvm.org/D86376 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
