[PATCH] D86376: [HIP] Simplify kernel launching

[Artem Belevich via Phabricator via cfe-commits]

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