llvm dev - Nov 2017 - CUDA kernel call with template parameters not appearing in Clang AST

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
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20171128/44ba592f/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/20171128/44ba592f/attachment-0001.png>

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-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.
>
>
> 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/279fc0cd/attachment.html>