zart Mo via llvm-dev
2021-Nov-21 14:27 UTC
[llvm-dev] 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. -------------- 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>