[PATCH] D44747: [AMDGPU] Set calling convention for CUDA kernel

John McCall via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Wed Mar 21 14:41:17 PDT 2018


rjmccall added a comment.

In https://reviews.llvm.org/D44747#1044916, @yaxunl wrote:

> In https://reviews.llvm.org/D44747#1044893, @rjmccall wrote:
>
> > Is there a reason for this to be done as a special case in IRGen instead of just implicitly applying the calling convention in Sema?
>
>
> The calling convention is not used in Sema, therefore it seems simpler to do it in codegen. I could try doing this in Sema too. Is there any advantage of doing this in Sema?


In IRGen, it's a special case for your specific language mode on your specific target.  In Sema, it can be done as part of the special checking for kernel functions.

Also, it looks like CUDA allows you to take the address of a __global__ function, and indirect calls to such functions presumably still follow the normal CUDA restrictions, so there must be *some* reflection of this in Sema.

Also, the calling convention


https://reviews.llvm.org/D44747





More information about the cfe-commits mailing list