[PATCH] D86376: [HIP] Simplify kernel launching

[Yaxun Liu via Phabricator via cfe-commits]

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