[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