[llvm-dev] CUDA kernel call with template parameters not appearing in Clang AST

[Michael Ehmke via llvm-dev]

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.

If I could get any assistance on why this is happening, I would appreciate it greatly.

Thank you,

Michael Ehmke

[Justin Lebar via llvm-dev]

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.

Regards,

-Justin

On Tue, Nov 28, 2017 at 2:33 PM Michael Ehmke via llvm-dev <llvm...@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.

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...@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

[Michael Ehmke via llvm-dev]

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 <jle...@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...@lists.llvm.org> wrote:

_______________________________________________

[Michael Ehmke via llvm-dev]

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