Michael Ehmke via llvm-dev
2017-Nov-29 14:05 UTC
[llvm-dev] CUDA kernel call with template parameters not appearing in Clang AST
Hi Justin, Thanks so much for the response. I'll try that out and get back to you. Sincerely, Michael Ehmke On Wed, Nov 29, 2017, 12:06 AM Justin Lebar <jlebar at google.com> wrote:> Hi, Michael. > > Are you able to create a reduced testcase which demonstrates your problem? > > I tried > > template <int N> > __global__ void Kernel() {} > > void test() { > Kernel<16><<<100, 200>>>(); > } > > which I compiled with > > $ clang -c -Xclang -ast-dump test.cu > > and the AST looks fine to me -- the CUDAKernelCallExpr shows up as > expected. > > -Justin > > On Tue, Nov 28, 2017 at 2:33 PM Michael Ehmke via llvm-dev < > llvm-dev at lists.llvm.org> wrote: > >> Hello all, >> >> I'm using Clang and am trying to refactor CUDA code. I want to traverse >> the AST to access a CUDAKernelCallExpr node where the kernel call has >> template parameters. I have been able to successfully match on and access >> CUDA kernel calls which have no template parameters by using the AST >> matchers to match on a CUDAKernelCallExpr. However, when my code comes >> across a kernel call with template parameters, it appears that the >> expression is not even existent in the AST when it is dumped. >> >> This is the expression that I'm having trouble parsing: >> >> ex. matrixMulCUDA<16><<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, >> dimsB.x); >> >> This expression does not show up at all in the AST when dumping the AST. >> I have been able to successfully parse the following kernel call expression >> and it shows up in the AST as a "CudaKernelCallExpr" node. >> >> ex. matrixMulCUDA<<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, dimsB.x); >> >> As you can see, the only difference is the <16> after the name of the >> call, and this node is clearly represented in the AST. >> >> Here is a screenshot showing the lack of a node to represent the call >> with template parameters. >> >> [image: image.png] >> If I could get any assistance on why this is happening, I would >> appreciate it greatly. >> >> Thank you, >> Michael Ehmke >> > _______________________________________________ >> LLVM Developers mailing list >> llvm-dev at lists.llvm.org >> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev >> >-------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20171129/26fd560b/attachment-0001.html> -------------- next part -------------- A non-text attachment was scrubbed... Name: image.png Type: image/png Size: 1368382 bytes Desc: not available URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20171129/26fd560b/attachment-0001.png>
Michael Ehmke via llvm-dev
2017-Nov-29 16:27 UTC
[llvm-dev] CUDA kernel call with template parameters not appearing in Clang AST
Hi Justin, To follow up, I tried out the sample code that you provided and the CUDAKernelCallExpr was in fact in the AST like you have said. I did some further troubleshooting with my code alongside making modifications to the code you provided me with and I discovered that the source of my issue was that I had been getting a missing file error for a header file for my source code (which had not interfered with anything for weeks now, so I disregarded it for the time-being) and the fact that I was getting that error was somehow interfering with the AST generation and causing the CUDAKernelCallExpr node to not appear in the AST only when it has template parameters. I added the includes directive for the missing header to your sample code and it had the same effect of causing the CUDAKernelCallExpr node to disappear. I have since resolved the missing dependency and all is well now with my original source code; all of the kernel call expressions are now showing up. Thank you for your help! Michael Ehmke On Wed, Nov 29, 2017 at 9:05 AM Michael Ehmke <mehmke at ncsu.edu> wrote:> Hi Justin, > > Thanks so much for the response. I'll try that out and get back to you. > > Sincerely, > Michael Ehmke > > > On Wed, Nov 29, 2017, 12:06 AM Justin Lebar <jlebar at google.com> wrote: > >> Hi, Michael. >> >> Are you able to create a reduced testcase which demonstrates your problem? >> >> I tried >> >> template <int N> >> __global__ void Kernel() {} >> >> void test() { >> Kernel<16><<<100, 200>>>(); >> } >> >> which I compiled with >> >> $ clang -c -Xclang -ast-dump test.cu >> >> and the AST looks fine to me -- the CUDAKernelCallExpr shows up as >> expected. >> >> -Justin >> >> On Tue, Nov 28, 2017 at 2:33 PM Michael Ehmke via llvm-dev < >> llvm-dev at lists.llvm.org> wrote: >> >>> Hello all, >>> >>> I'm using Clang and am trying to refactor CUDA code. I want to traverse >>> the AST to access a CUDAKernelCallExpr node where the kernel call has >>> template parameters. I have been able to successfully match on and access >>> CUDA kernel calls which have no template parameters by using the AST >>> matchers to match on a CUDAKernelCallExpr. However, when my code comes >>> across a kernel call with template parameters, it appears that the >>> expression is not even existent in the AST when it is dumped. >>> >>> This is the expression that I'm having trouble parsing: >>> >>> ex. matrixMulCUDA<16><<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, >>> dimsB.x); >>> >>> This expression does not show up at all in the AST when dumping the AST. >>> I have been able to successfully parse the following kernel call expression >>> and it shows up in the AST as a "CudaKernelCallExpr" node. >>> >>> ex. matrixMulCUDA<<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, >>> dimsB.x); >>> >>> As you can see, the only difference is the <16> after the name of the >>> call, and this node is clearly represented in the AST. >>> >>> Here is a screenshot showing the lack of a node to represent the call >>> with template parameters. >>> >>> [image: image.png] >>> If I could get any assistance on why this is happening, I would >>> appreciate it greatly. >>> >>> Thank you, >>> Michael Ehmke >>> >> _______________________________________________ >>> LLVM Developers mailing list >>> llvm-dev at lists.llvm.org >>> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev >>> >>-------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20171129/3e85bac9/attachment-0001.html> -------------- next part -------------- A non-text attachment was scrubbed... Name: image.png Type: image/png Size: 1368382 bytes Desc: not available URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20171129/3e85bac9/attachment-0001.png>