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

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.

sample.ll (53.4 KB)