yaxunl added a comment. In D86376#2239618 <https://reviews.llvm.org/D86376#2239618>, @tra wrote:
> 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? Yes that should work. Will do. 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
