[llvm-dev] How to get the passed Argument of every GPU kernel CallInst in LLVM-IR produced by HIPCC

[zart Mo via llvm-dev]

I am developing a data-flow analysis optimization pass on the IR produced
from a HIP program. But when I am trying to get the input Value of a
CallInst calling GPU kernel like what we used to do on a CallInst of CPU
function, I found that it's totally different.

First, all GPU kernel CallInst is like:

> %%105 = call i32 @hipLaunchKernel(i8* bitcast (void (float*, float*, float*,
> i32, i32)* @_Z30__device_stub__vectoradd_floatPfPKfS1_ii to i8*), i64 %98,
> i32 %100, i64 %102, i32 %104, i8** nonnull %83, i64 %94,
> %struct.ihipStream_t* %96
>

I cannot connect them with the variable which I defined to be used for the
passed argument in HIP program via use-def chain.

Second, when I look further, some instructions producing these arguments
which are involved in this callinst, has a context called !13, and they are
like something below:

> %98 = load i64, i64* %97, align 8, !noalias !13
>
which locates in:

> !13 = !{!14, !16, !17}
> !14 = distinct !{!14, !15, !"_Z30__device_stub__vectoradd_floatPfPKfS1_ii:
> argument 0"}
> !15 = distinct !{!15, !"_Z30__device_stub__vectoradd_floatPfPKfS1_ii"}
> !16 = distinct !{!16, !15, !"_Z30__device_stub__vectoradd_floatPfPKfS1_ii:
> argument 1"}
> !17 = distinct !{!17, !15, !"_Z30__device_stub__vectoradd_floatPfPKfS1_ii:
> argument 2"}
>

So now I am confused about how to track the input Value of each GPU kernel
CallInst in such a situation.
If I have to take them through the !13 variable, there are some IR files
that are not including such variables to keep arguments of gpu kernel in
them, and what should I do about those? Is there some specific rules for
handling such a IR file produced from HIP program?

I also posted my IR file for people who are interested in this. I really
appreciate any help or discussion on this, thank you.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20211121/2ea0951e/attachment-0001.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: sample.ll
Type: application/octet-stream
Size: 54650 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20211121/2ea0951e/attachment-0001.obj>