[PATCH] D44747: Set calling convention for CUDA kernel

[Artem Belevich via Phabricator via llvm-commits]

tra added a comment.

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

> I will try fixing that.
>
> The CUDA kernel calling convention should be dropped in all DRE's since it is invisible to the user.
>
> Sam


Apparently it's not always the case.
Here's a bit more elaborate example demonstrating inconsistency in this behavior. Calling convention is ignored for regular functions, but not for function templates.

  __global__ void EmptyKernel(void) { }
  
  template <typename T>
  __global__ void EmptyKernelT(void) { }
  
  struct Dummy {
    /// Type definition of the EmptyKernel kernel entry point
    typedef void (*EmptyKernelPtr)();
    EmptyKernelPtr Empty() { return EmptyKernel; } // this one works
    EmptyKernelPtr EmptyT() { return EmptyKernelT<void>; } // this one errors out.
  };

Do you think this is something you can fix quickly or do you want to unroll the change while you figure out what's going on?


Repository:
  rL LLVM

https://reviews.llvm.org/D44747